From 8928835f6b87b618eccdfcd4b88495fe8f5d1a10 Mon Sep 17 00:00:00 2001 From: Raphael Amorim Date: Wed, 29 Apr 2026 02:03:27 +0200 Subject: [PATCH] fix rendering order for images --- sugarloaf/src/grid/cpu.rs | 43 ++++-- sugarloaf/src/grid/metal.rs | 258 ++++++++++++++++++++++------------ sugarloaf/src/grid/mod.rs | 102 ++++++++++---- sugarloaf/src/grid/vulkan.rs | 33 +++-- sugarloaf/src/grid/webgpu.rs | 28 ++-- sugarloaf/src/renderer/cpu.rs | 11 +- sugarloaf/src/renderer/mod.rs | 140 ++++++++++++++---- sugarloaf/src/sugarloaf.rs | 22 ++- 8 files changed, 455 insertions(+), 182 deletions(-) diff --git a/sugarloaf/src/grid/cpu.rs b/sugarloaf/src/grid/cpu.rs index 6bb7991c..beec1ecd 100644 --- a/sugarloaf/src/grid/cpu.rs +++ b/sugarloaf/src/grid/cpu.rs @@ -340,7 +340,9 @@ impl CpuGridRenderer { /// caller's `0x00RRGGBB` u32 buffer. Mirrors the bg + text passes /// of `grid.metal`'s shaders, in the same draw order so glyphs /// composite correctly over their cell backgrounds. - pub fn render( + /// Paint the cell-bg pass into `buf`. Pair with `render_text`, + /// with any `kitty_below_text` images composited in between. + pub fn render_bg( &self, buf: &mut [u32], buf_w: u32, @@ -357,7 +359,6 @@ impl CpuGridRenderer { if cols == 0 || rows == 0 { return; } - // grid_padding = (top, right, bottom, left) — same as the shader. let pad_top = uniforms.grid_padding[0]; let pad_left = uniforms.grid_padding[3]; @@ -366,11 +367,8 @@ impl CpuGridRenderer { let cursor_x = uniforms.cursor_pos[0]; let cursor_y = uniforms.cursor_pos[1]; let cursor_bg_active = uniforms.cursor_bg_color[3] > 0.0; - let cursor_fg_active = uniforms.cursor_color[3] > 0.0; let cursor_bg = normalize_color(uniforms.cursor_bg_color); - let cursor_fg = normalize_color(uniforms.cursor_color); - // ---------- bg pass ---------- let buf_cols = self.cols as usize; let row_count = (rows as usize).min(self.rows as usize); let col_count = (cols as usize).min(self.cols as usize); @@ -392,8 +390,38 @@ impl CpuGridRenderer { fill_rect(buf, buf_w_i, buf_h_i, x0, y0, x1, y1, rgba); } } + } + + /// Paint the cell-text pass into `buf`. Walks `fg_rows` and blits + /// each glyph's atlas slot at the cell origin computed from + /// `grid_pos + bearings`. + pub fn render_text( + &self, + buf: &mut [u32], + buf_w: u32, + buf_h: u32, + uniforms: &GridUniforms, + ) { + let cell_w = uniforms.cell_size[0]; + let cell_h = uniforms.cell_size[1]; + if cell_w <= 0.0 || cell_h <= 0.0 { + return; + } + let cols = uniforms.grid_size[0]; + let rows = uniforms.grid_size[1]; + if cols == 0 || rows == 0 { + return; + } + let pad_top = uniforms.grid_padding[0]; + let pad_left = uniforms.grid_padding[3]; + + let buf_w_i = buf_w as i32; + let buf_h_i = buf_h as i32; + let cursor_x = uniforms.cursor_pos[0]; + let cursor_y = uniforms.cursor_pos[1]; + let cursor_fg_active = uniforms.cursor_color[3] > 0.0; + let cursor_fg = normalize_color(uniforms.cursor_color); - // ---------- text pass ---------- let mask = self.atlas_grayscale.pixels(); let mask_side = self.atlas_grayscale.side as usize; let color_atlas = self.atlas_color.pixels(); @@ -407,11 +435,8 @@ impl CpuGridRenderer { continue; } - // Cell origin in pixels (matches grid.metal:258 + 275-276). let cell_pos_x = (glyph.grid_pos[0] as f32) * cell_w + pad_left; let cell_pos_y = (glyph.grid_pos[1] as f32) * cell_h + pad_top; - // Glyph origin (matches grid.metal:267-271). bearings.y is - // measured from the cell bottom, so flip into top-down space. let glyph_x = (cell_pos_x + glyph.bearings[0] as f32) as i32; let glyph_y = (cell_pos_y + cell_h - glyph.bearings[1] as f32) as i32; diff --git a/sugarloaf/src/grid/metal.rs b/sugarloaf/src/grid/metal.rs index 7e8d65c0..356ce9c3 100644 --- a/sugarloaf/src/grid/metal.rs +++ b/sugarloaf/src/grid/metal.rs @@ -23,17 +23,66 @@ use metal::{ RenderPipelineDescriptor, RenderPipelineState, Texture, TextureDescriptor, VertexDescriptor, }; +use parking_lot::{Condvar, Mutex}; use rustc_hash::FxHashMap; +use std::sync::Arc; use super::atlas::{AtlasSlot, GlyphKey, RasterizedGlyph}; use super::cell::{CellBg, CellText, GridUniforms}; use crate::context::metal::MetalContext; use crate::renderer::image_cache::atlas::AtlasAllocator; -/// Reserved for Phase 1c's completion-handler-gated triple buffering. -/// Currently only slot 0 is used. +/// Number of GPU buffer slots cycled through frame-by-frame so the +/// CPU can populate frame N+1 while the GPU is still drawing N. Sized +/// to match Metal `swap_chain_count` at +/// `ghostty/src/renderer/Metal.zig:37`. const FRAMES_IN_FLIGHT: usize = 3; +/// Number of in-flight frames the swap chain owns. Re-export here +/// so the Renderer-level swap chain (which advances frame_index + +/// owns the permit pool) can size itself off the same constant the +/// per-grid buffer arrays use. +pub const FRAMES_IN_FLIGHT_PUB: usize = FRAMES_IN_FLIGHT; + +/// Cross-thread semaphore — `FRAMES_IN_FLIGHT_PUB` permits, +/// decremented on render-start, restored in the command-buffer +/// completion handler. Mirrors ghostty's `frame_sema` at +/// `renderer/generic.zig:261` (`std.Thread.Semaphore` → +/// `dispatch_semaphore_t` on Darwin). We use parking_lot's +/// Mutex+Condvar instead of dispatch_semaphore because the rest of +/// sugarloaf already pulls parking_lot in and the behavior is +/// identical for a counting semaphore. +/// +/// One pool is shared across all grids on a `Renderer`, so a render +/// only acquires one permit even with N split panels. +pub type FramePermits = Arc<(Mutex, Condvar)>; + +/// Construct a permit pool with `FRAMES_IN_FLIGHT` initial permits. +pub fn new_frame_permits() -> FramePermits { + Arc::new((Mutex::new(FRAMES_IN_FLIGHT), Condvar::new())) +} + +/// Block until a permit is available, then take one. Called once per +/// frame at render start. +pub fn acquire_frame_permit(p: &FramePermits) { + let (m, c) = &**p; + let mut g = m.lock(); + while *g == 0 { + c.wait(&mut g); + } + *g -= 1; +} + +/// Release a permit, signaling any frame waiting on +/// `acquire_frame_permit`. Called from the command-buffer completion +/// handler — i.e. on a Metal-internal thread, not the main thread. +pub fn release_frame_permit(p: &FramePermits) { + let (m, c) = &**p; + let mut g = m.lock(); + *g += 1; + c.notify_one(); +} + /// Extra slots appended to the per-row fg storage for cursor glyphs. /// `rows + 2` layout (block cursor at slot 0, /// non-block-style cursor at the tail). @@ -250,57 +299,60 @@ pub struct MetalGridRenderer { cols: u32, rows: u32, - /// `cols * rows` CellBg entries. Triple-buffered storage is - /// allocated for Phase 1c; only `bg_buffers[0]` is active in - /// Phase 1a. + /// CPU-side shadow of the per-cell bg buffer. `write_row` / + /// `clear_row` mutate this directly; `render_bg` flushes it into + /// the active frame slot's GPU buffer when the slot is dirty. + /// Decoupling the writes from the GPU buffer is what makes + /// triple-buffering safe — the CPU can mutate `bg_cpu` while the + /// GPU is still reading from a previous slot. + bg_cpu: Vec, + + /// One GPU `CellBg` buffer per in-flight frame slot. The slot + /// active for a given frame is selected by `next_frame()` and + /// stored in `self.frame`. bg_buffers: [Buffer; FRAMES_IN_FLIGHT], - /// Per-row FG glyph storage. Slot 0 = block cursor cells, - /// 1..=rows = content rows, last = non-block cursor cells. Unused - /// in Phase 1a (Phase 1c turns this on alongside the text shader). - #[allow(dead_code)] + /// Per-slot dirty flag. Any `write_row` / `clear_row` sets all + /// `FRAMES_IN_FLIGHT` flags so each slot re-flushes from `bg_cpu` + /// when it next becomes the active frame; the slot's flag is + /// cleared after the per-slot flush in `render_bg`. + bg_dirty: [bool; FRAMES_IN_FLIGHT], + + /// Per-row FG glyph storage (CPU-side, single copy across all + /// frame slots). Slot 0 = block cursor cells, 1..=rows = content + /// rows, last = non-block cursor cells. fg_rows: Vec>, - /// GPU buffer that holds the concatenation of all `fg_rows`. - /// Reserved for Phase 1c. - #[allow(dead_code)] + /// One per-instance `CellText` vertex buffer per in-flight frame + /// slot. Re-uploaded from `fg_rows` on the slot's first frame + /// after a write. fg_buffers: [Buffer; FRAMES_IN_FLIGHT], - #[allow(dead_code)] fg_capacity: [usize; FRAMES_IN_FLIGHT], - /// Ring index — Phase 1a always reads/writes 0. - #[allow(dead_code)] - frame: usize, - /// Compiled bg render pipeline. Binds: /// buffer(0): `GridUniforms` (via `set_vertex_bytes` / /// `set_fragment_bytes`) - /// buffer(1): `bg_buffers[0]` + /// buffer(1): `bg_buffers[frame]` bg_pipeline: RenderPipelineState, /// Compiled text render pipeline. Binds: /// buffer(0): per-instance `CellText` vertex buffer /// buffer(1): `GridUniforms` /// texture(0): `atlas_grayscale` - /// texture(1): `atlas_color` (reused = atlas_grayscale for now) + /// texture(1): `atlas_color` text_pipeline: RenderPipelineState, /// Staging buffer for the concatenated fg instances. Rebuilt each /// frame by flattening `fg_rows` into a contiguous slice. fg_staging: Vec, - /// Instance count that's live on the GPU in `fg_buffers[0]` from - /// the previous render. When no row has been touched since, the - /// GPU copy is already correct and we skip the concat + upload, - /// reissuing the same `draw_primitives_instanced` call against - /// the resident data. Invalidated by `write_row` / `clear_row` / - /// `resize`. - fg_live_count: u32, + /// Per-slot live instance count from the most recent flush. + fg_live_count: [u32; FRAMES_IN_FLIGHT], - /// `true` when `fg_rows` holds writes not yet flushed to - /// `fg_buffers`. Set by any row-level write, cleared after a - /// successful flush in `render`. - fg_dirty: bool, + /// Per-slot dirty flag for the fg path. Any row-level write sets + /// all flags; the slot's flag is cleared by `render_text` after a + /// successful flush. + fg_dirty: [bool; FRAMES_IN_FLIGHT], /// Grayscale (R8) glyph atlas — outline mask bitmaps from the /// monochrome rasterizer path. @@ -337,21 +389,25 @@ impl MetalGridRenderer { let atlas_grayscale = MetalGlyphAtlas::new_grayscale(&device); let atlas_color = MetalGlyphAtlas::new_color(&device); + let bg_cpu_len = (cols as usize) * (rows as usize); + let bg_cpu = vec![CellBg::TRANSPARENT; bg_cpu_len]; + Self { device, command_queue, cols, rows, + bg_cpu, bg_buffers, + bg_dirty: [true; FRAMES_IN_FLIGHT], fg_rows: init_fg_rows(rows), fg_buffers, fg_capacity, - frame: 0, bg_pipeline, text_pipeline, fg_staging: Vec::new(), - fg_live_count: 0, - fg_dirty: true, + fg_live_count: [0; FRAMES_IN_FLIGHT], + fg_dirty: [true; FRAMES_IN_FLIGHT], atlas_grayscale, atlas_color, needs_full_rebuild: true, @@ -420,6 +476,7 @@ impl MetalGridRenderer { } self.cols = cols; self.rows = rows; + self.bg_cpu = vec![CellBg::TRANSPARENT; (cols as usize) * (rows as usize)]; self.bg_buffers = std::array::from_fn(|_| alloc_bg_buffer(&self.device, cols, rows)); self.fg_rows = init_fg_rows(rows); @@ -428,20 +485,24 @@ impl MetalGridRenderer { std::array::from_fn(|_| alloc_fg_buffer(&self.device, initial_fg_capacity)); self.fg_capacity = [initial_fg_capacity; FRAMES_IN_FLIGHT]; // Fresh buffers = zero contents; emission path must rewrite - // every row on the next frame even if no damage came in. + // every row on the next frame even if no damage came in. All + // 3 in-flight slots are now stale so each one needs a flush + // when it next becomes the active frame. self.needs_full_rebuild = true; - self.fg_dirty = true; - self.fg_live_count = 0; + self.bg_dirty = [true; FRAMES_IN_FLIGHT]; + self.fg_dirty = [true; FRAMES_IN_FLIGHT]; + self.fg_live_count = [0; FRAMES_IN_FLIGHT]; } pub fn write_row(&mut self, row: u32, bg: &[CellBg], fg: &[CellText]) { - // FG: stash in the CPU-side per-row vec. Phase 1c will - // concatenate these into a GPU buffer at render time. + // FG: stash in the CPU-side per-row vec. All in-flight slots + // need re-flushing — set every dirty flag so each slot + // re-uploads when it becomes the active frame. let idx = (row as usize) + 1; if let Some(slot) = self.fg_rows.get_mut(idx) { slot.clear(); slot.extend_from_slice(fg); - self.fg_dirty = true; + self.fg_dirty = [true; FRAMES_IN_FLIGHT]; } if row >= self.rows { @@ -449,23 +510,19 @@ impl MetalGridRenderer { } let row_start = (row as usize) * (self.cols as usize); let row_len = (self.cols as usize).min(bg.len()); - let buf = &self.bg_buffers[0]; - unsafe { - let ptr = buf.contents() as *mut CellBg; - let dst = - std::slice::from_raw_parts_mut(ptr.add(row_start), self.cols as usize); - dst[..row_len].copy_from_slice(&bg[..row_len]); - for slot in &mut dst[row_len..] { - *slot = CellBg::TRANSPARENT; - } + let dst = &mut self.bg_cpu[row_start..row_start + self.cols as usize]; + dst[..row_len].copy_from_slice(&bg[..row_len]); + for slot in &mut dst[row_len..] { + *slot = CellBg::TRANSPARENT; } + self.bg_dirty = [true; FRAMES_IN_FLIGHT]; } pub fn clear_row(&mut self, row: u32) { let idx = (row as usize) + 1; if let Some(slot) = self.fg_rows.get_mut(idx) { if !slot.is_empty() { - self.fg_dirty = true; + self.fg_dirty = [true; FRAMES_IN_FLIGHT]; } slot.clear(); } @@ -473,14 +530,13 @@ impl MetalGridRenderer { return; } let row_start = (row as usize) * (self.cols as usize); - let buf = &self.bg_buffers[0]; - unsafe { - let ptr = buf.contents() as *mut CellBg; - let dst = - std::slice::from_raw_parts_mut(ptr.add(row_start), self.cols as usize); - for slot in dst { - *slot = CellBg::TRANSPARENT; - } + let dst = &mut self.bg_cpu[row_start..row_start + self.cols as usize]; + let needs_flush = dst.iter().any(|c| *c != CellBg::TRANSPARENT); + for slot in dst { + *slot = CellBg::TRANSPARENT; + } + if needs_flush { + self.bg_dirty = [true; FRAMES_IN_FLIGHT]; } } @@ -495,7 +551,7 @@ impl MetalGridRenderer { } slot.clear(); slot.extend_from_slice(cells); - self.fg_dirty = true; + self.fg_dirty = [true; FRAMES_IN_FLIGHT]; } } @@ -512,7 +568,7 @@ impl MetalGridRenderer { } slot.clear(); slot.extend_from_slice(cells); - self.fg_dirty = true; + self.fg_dirty = [true; FRAMES_IN_FLIGHT]; } } @@ -538,20 +594,39 @@ impl MetalGridRenderer { } } if changed { - self.fg_dirty = true; + self.fg_dirty = [true; FRAMES_IN_FLIGHT]; } } - /// Record both grid passes against the caller's `encoder`. The - /// caller owns the command buffer, drawable, and render pass - /// descriptor. Draw order: - /// - /// 1. bg pass — fullscreen triangle, per-fragment cell lookup. - /// 2. text pass — one instanced quad per `CellText` in `fg_rows`. - pub fn render(&mut self, encoder: &RenderCommandEncoderRef, uniforms: &GridUniforms) { - let uniforms_bytes = bytemuck::bytes_of(uniforms); + /// Record the cell-bg pass against the caller's `encoder`. Drawn + /// as a fullscreen triangle; the fragment shader resolves each + /// pixel back to its owning cell and reads the per-cell color + /// from `bg_buffers[frame]`. Caller owns command buffer + + /// drawable + pass descriptor; `frame` comes from the Renderer's + /// shared swap chain (acquired via `acquire_frame_permit` once + /// per render). Pair with `render_text` after compositing any + /// `kitty_below_text` images in between (matches + /// `renderer/generic.zig:1654-1668`). + pub fn render_bg( + &mut self, + encoder: &RenderCommandEncoderRef, + frame: usize, + uniforms: &GridUniforms, + ) { + // Flush the CPU shadow into this slot's GPU buffer if it + // hasn't been flushed since the last `write_row`. Each slot + // tracks its own dirty bit so all 3 in-flight slots stay + // consistent without re-uploading on every frame. + if self.bg_dirty[frame] { + let bytes = bytemuck::cast_slice::(&self.bg_cpu); + unsafe { + let dst = self.bg_buffers[frame].contents() as *mut u8; + std::ptr::copy_nonoverlapping(bytes.as_ptr(), dst, bytes.len()); + } + self.bg_dirty[frame] = false; + } - // ---------- bg pass ---------- + let uniforms_bytes = bytemuck::bytes_of(uniforms); encoder.set_render_pipeline_state(&self.bg_pipeline); encoder.set_vertex_bytes( 0, @@ -563,17 +638,23 @@ impl MetalGridRenderer { uniforms_bytes.len() as u64, uniforms_bytes.as_ptr() as *const std::ffi::c_void, ); - encoder.set_fragment_buffer(1, Some(&self.bg_buffers[0]), 0); + encoder.set_fragment_buffer(1, Some(&self.bg_buffers[frame]), 0); encoder.draw_primitives(MTLPrimitiveType::Triangle, 0, 3); + } - // ---------- text pass ---------- - // When no row has been written since the last flush, `fg_buffers[0]` - // already holds the exact same instances the GPU needs — skip the - // concat + memcpy and just re-bind + re-draw. On a Noop/CursorOnly - // damage frame (blink tick, scrollbar fade, etc.) this is the - // whole difference between ~0 µs and ~(rows × cols × 32 B) of - // wasted CPU work per frame. - if self.fg_dirty { + /// Record the cell-text pass. One instanced quad per `CellText` + /// in `fg_rows`. Lazily flushes the per-row CPU vecs into + /// `fg_buffers[frame]` only when the slot is dirty — on a + /// Noop/CursorOnly damage frame (blink tick, scrollbar fade, + /// etc.) this is the whole difference between ~0 µs and ~(rows × + /// cols × 32 B) of wasted CPU work per frame. + pub fn render_text( + &mut self, + encoder: &RenderCommandEncoderRef, + frame: usize, + uniforms: &GridUniforms, + ) { + if self.fg_dirty[frame] { // Flatten per-row fg_rows into the staging vec. Order matters // for z: slot 0 (block cursor) first, content rows next, // non-block-cursor slot last — same approach's ordering. @@ -582,34 +663,32 @@ impl MetalGridRenderer { self.fg_staging.extend_from_slice(row); } - // Grow the GPU buffer if the staging vec outran current capacity. - if self.fg_staging.len() > self.fg_capacity[0] { + if self.fg_staging.len() > self.fg_capacity[frame] { let new_cap = self.fg_staging.len().next_power_of_two(); - self.fg_buffers[0] = alloc_fg_buffer(&self.device, new_cap); - self.fg_capacity[0] = new_cap; + self.fg_buffers[frame] = alloc_fg_buffer(&self.device, new_cap); + self.fg_capacity[frame] = new_cap; } // Upload staging → GPU buffer. Shared storage mode means the // CPU pointer is the GPU pointer. let fg_bytes = bytemuck::cast_slice::(&self.fg_staging); unsafe { - let dst = self.fg_buffers[0].contents() as *mut u8; + let dst = self.fg_buffers[frame].contents() as *mut u8; std::ptr::copy_nonoverlapping(fg_bytes.as_ptr(), dst, fg_bytes.len()); } - self.fg_live_count = self.fg_staging.len() as u32; - self.fg_dirty = false; + self.fg_live_count[frame] = self.fg_staging.len() as u32; + self.fg_dirty[frame] = false; } - let instance_count = self.fg_live_count as usize; + let instance_count = self.fg_live_count[frame] as usize; if instance_count == 0 { return; } + let uniforms_bytes = bytemuck::bytes_of(uniforms); encoder.set_render_pipeline_state(&self.text_pipeline); - // buffer(0): per-instance vertex data. - encoder.set_vertex_buffer(0, Some(&self.fg_buffers[0]), 0); - // buffer(1): uniforms (pushed inline). + encoder.set_vertex_buffer(0, Some(&self.fg_buffers[frame]), 0); encoder.set_vertex_bytes( 1, uniforms_bytes.len() as u64, @@ -618,7 +697,6 @@ impl MetalGridRenderer { encoder.set_fragment_texture(0, Some(&self.atlas_grayscale.texture)); encoder.set_fragment_texture(1, Some(&self.atlas_color.texture)); - // Four-vertex triangle strip per instance (the quad). encoder.draw_primitives_instanced( MTLPrimitiveType::TriangleStrip, 0, diff --git a/sugarloaf/src/grid/mod.rs b/sugarloaf/src/grid/mod.rs index 5be84666..081517f3 100644 --- a/sugarloaf/src/grid/mod.rs +++ b/sugarloaf/src/grid/mod.rs @@ -180,41 +180,62 @@ impl GridRenderer { } } - /// Record grid draw calls against a caller-supplied render pass / - /// encoder. The caller owns the command buffer + drawable + pass - /// descriptor so the grid composes with sugarloaf's UI overlays - /// (island, assistant, etc.) in a single render pass. - /// - /// Phase 1a: Metal draws the bg pass; Wgpu is still a no-op. + /// Record the cell-bg pass for this grid. `frame` is the swap + /// chain slot index acquired from the Renderer's shared + /// `FramePermits` pool — the grid uses it to pick the right + /// per-slot GPU buffer. Caller composites `kitty_below_text` + /// images between this call and `render_text_*` to match + /// `renderer/generic.zig:1654-1668` ordering. #[cfg(target_os = "macos")] - pub fn render_metal( + pub fn render_bg_metal( &mut self, encoder: &::metal::RenderCommandEncoderRef, + frame: usize, uniforms: &GridUniforms, ) { if let GridRenderer::Metal(r) = self { - r.render(encoder, uniforms); + r.render_bg(encoder, frame, uniforms); } } - /// Wgpu counterpart of `render_metal`. Phase 1b will record a bg - /// pass against the caller's `wgpu::RenderPass`. + /// Record the cell-text pass for this grid. Caller composites + /// `kitty_above_text` images after this call. + #[cfg(target_os = "macos")] + pub fn render_text_metal( + &mut self, + encoder: &::metal::RenderCommandEncoderRef, + frame: usize, + uniforms: &GridUniforms, + ) { + if let GridRenderer::Metal(r) = self { + r.render_text(encoder, frame, uniforms); + } + } + + /// Wgpu cell-bg pass. Pair with `render_text_wgpu`. + #[cfg(feature = "wgpu")] + pub fn render_bg_wgpu( + &mut self, + render_pass: &mut wgpu::RenderPass<'_>, + uniforms: &GridUniforms, + ) { + if let GridRenderer::Wgpu(r) = self { + r.render_bg(render_pass, uniforms); + } + } + + /// Wgpu cell-text pass. #[cfg(feature = "wgpu")] - pub fn render_wgpu<'pass>( - &'pass mut self, - render_pass: &mut wgpu::RenderPass<'pass>, + pub fn render_text_wgpu( + &mut self, + render_pass: &mut wgpu::RenderPass<'_>, uniforms: &GridUniforms, ) { if let GridRenderer::Wgpu(r) = self { - r.render(render_pass, uniforms); + r.render_text(render_pass, uniforms); } } - /// Vulkan counterpart of `render_metal`. Records draws into - /// `cmd_buffer`, which the caller (`Sugarloaf::render_vulkan`) - /// has already wrapped in `cmd_begin_rendering` against the - /// swapchain image. No-op when this renderer isn't a Vulkan - /// variant — same shape as `render_wgpu` / `render_metal`. /// Pre-pass hook: flush atlas uploads before the caller opens /// dynamic rendering. Must be called BEFORE `cmd_begin_rendering` /// because `vkCmdCopyBufferToImage` is forbidden inside a render @@ -233,8 +254,9 @@ impl GridRenderer { } } + /// Vulkan cell-bg pass. #[cfg(target_os = "linux")] - pub fn render_vulkan( + pub fn render_bg_vulkan( &mut self, ctx: &crate::context::vulkan::VulkanContext, cmd_buffer: ash::vk::CommandBuffer, @@ -242,15 +264,41 @@ impl GridRenderer { uniforms: &GridUniforms, ) { if let GridRenderer::Vulkan(r) = self { - r.render(ctx, cmd_buffer, frame_slot, uniforms); + r.render_bg(ctx, cmd_buffer, frame_slot, uniforms); + } + } + + /// Vulkan cell-text pass. + #[cfg(target_os = "linux")] + pub fn render_text_vulkan( + &mut self, + ctx: &crate::context::vulkan::VulkanContext, + cmd_buffer: ash::vk::CommandBuffer, + frame_slot: usize, + uniforms: &GridUniforms, + ) { + if let GridRenderer::Vulkan(r) = self { + r.render_text(ctx, cmd_buffer, frame_slot, uniforms); + } + } + + /// Software cell-bg pass. Paints the grid bg into the + /// caller-supplied `0x00RRGGBB` u32 buffer (typically + /// softbuffer's `buffer_mut`). No-op for non-CPU variants. + pub fn render_bg_cpu( + &self, + buf: &mut [u32], + buf_w: u32, + buf_h: u32, + uniforms: &GridUniforms, + ) { + if let GridRenderer::Cpu(r) = self { + r.render_bg(buf, buf_w, buf_h, uniforms); } } - /// Software counterpart of `render_metal` / `render_vulkan` / - /// `render_wgpu`. Paints the grid into the caller-supplied - /// `0x00RRGGBB` u32 buffer (typically softbuffer's `buffer_mut`). - /// No-op for non-CPU variants. - pub fn render_cpu( + /// Software cell-text pass. + pub fn render_text_cpu( &self, buf: &mut [u32], buf_w: u32, @@ -258,7 +306,7 @@ impl GridRenderer { uniforms: &GridUniforms, ) { if let GridRenderer::Cpu(r) = self { - r.render(buf, buf_w, buf_h, uniforms); + r.render_text(buf, buf_w, buf_h, uniforms); } } diff --git a/sugarloaf/src/grid/vulkan.rs b/sugarloaf/src/grid/vulkan.rs index e880d3ba..5bd7077a 100644 --- a/sugarloaf/src/grid/vulkan.rs +++ b/sugarloaf/src/grid/vulkan.rs @@ -612,14 +612,15 @@ impl VulkanGridRenderer { self.flush_pending_uploads(ctx, cmd, frame_slot); } - /// Record the bg + text passes into `cmd`. Caller has already - /// opened the dynamic-rendering pass and set viewport/scissor. - /// `frame_slot` is the in-flight slot whose `in_flight` fence has - /// been waited on. Atlas uploads must already have been flushed - /// via `prepare()` before the pass opened. - pub fn render( + /// Record the cell-bg pass into `cmd`. Uploads bg cells + + /// uniforms for `frame_slot` first, then issues the fullscreen + /// triangle. Caller must have opened the dynamic-rendering pass + + /// set viewport/scissor + flushed atlas uploads via `prepare()`. + /// Pair with `render_text`, with any `kitty_below_text` images + /// composited in between. + pub fn render_bg( &mut self, - ctx: &VulkanContext, + _ctx: &VulkanContext, cmd: vk::CommandBuffer, frame_slot: usize, uniforms: &GridUniforms, @@ -627,7 +628,6 @@ impl VulkanGridRenderer { debug_assert!(frame_slot < FRAMES_IN_FLIGHT); let slot = frame_slot; - // ----- bg cells + uniforms upload ----- if self.bg_dirty[slot] { unsafe { let dst = self.bg_buffers[slot].as_mut_ptr() as *mut CellBg; @@ -644,7 +644,6 @@ impl VulkanGridRenderer { std::ptr::write(dst, *uniforms); } - // ----- bg pass (1 fullscreen triangle, fragment does cell lookup) ----- unsafe { self.device.cmd_bind_pipeline( cmd, @@ -661,8 +660,22 @@ impl VulkanGridRenderer { ); self.device.cmd_draw(cmd, 3, 1, 0, 0); } + } + + /// Record the cell-text pass into `cmd`. One instanced quad per + /// `CellText`. Lazily flushes per-row CPU vecs into the per-slot + /// fg buffer only when `fg_dirty[slot]` — no concat + memcpy on + /// Noop/CursorOnly damage frames. + pub fn render_text( + &mut self, + ctx: &VulkanContext, + cmd: vk::CommandBuffer, + frame_slot: usize, + _uniforms: &GridUniforms, + ) { + debug_assert!(frame_slot < FRAMES_IN_FLIGHT); + let slot = frame_slot; - // ----- text pass (instanced quads, one per glyph) ----- if self.fg_dirty[slot] { self.fg_staging.clear(); for row in &self.fg_rows { diff --git a/sugarloaf/src/grid/webgpu.rs b/sugarloaf/src/grid/webgpu.rs index 862fe9fb..08fb1d68 100644 --- a/sugarloaf/src/grid/webgpu.rs +++ b/sugarloaf/src/grid/webgpu.rs @@ -440,19 +440,19 @@ impl WgpuGridRenderer { self.atlas_color.insert(key, glyph) } - /// Record bg pass + text pass against the caller's `render_pass`. - pub fn render<'pass>( - &'pass mut self, - render_pass: &mut wgpu::RenderPass<'pass>, + /// Record the cell-bg pass. Uploads the uniform buffer (cheap; the + /// bg path always runs first per frame so this is the right place + /// for it) and the bg cell storage buffer if it changed since the + /// last flush. Pair with `render_text`, with any + /// `kitty_below_text` images composited in between. + pub fn render_bg( + &mut self, + render_pass: &mut wgpu::RenderPass<'_>, uniforms: &GridUniforms, ) { - // Uniforms always upload (cheap, and cursor/min_contrast can - // change without a row write). self.queue .write_buffer(&self.uniform_buffer, 0, bytemuck::bytes_of(uniforms)); - // Skip re-uploading bg cells when no row changed — the GPU - // copy is already correct from the previous frame. if self.bg_dirty { self.queue.write_buffer( &self.bg_buffers[0], @@ -462,12 +462,19 @@ impl WgpuGridRenderer { self.bg_dirty = false; } - // ---------- bg pass ---------- render_pass.set_pipeline(&self.bg_pipeline); render_pass.set_bind_group(0, &self.bg_bind_group, &[]); render_pass.draw(0..3, 0..1); + } - // ---------- text pass ---------- + /// Record the cell-text pass. Lazily flushes per-row CPU vecs to + /// `fg_buffers[0]` only when `fg_dirty` — saves ~(rows × cols × + /// 32 B) of memcpy on Noop/CursorOnly damage frames. + pub fn render_text( + &mut self, + render_pass: &mut wgpu::RenderPass<'_>, + _uniforms: &GridUniforms, + ) { if self.fg_dirty { self.fg_staging.clear(); for row in &self.fg_rows { @@ -497,7 +504,6 @@ impl WgpuGridRenderer { render_pass.set_bind_group(0, &self.text_uniform_bg, &[]); render_pass.set_bind_group(1, &self.text_atlas_bg, &[]); render_pass.set_vertex_buffer(0, self.fg_buffers[0].slice(..)); - // 4 vertices per instance → triangle strip quad. render_pass.draw(0..4, 0..instance_count as u32); } } diff --git a/sugarloaf/src/renderer/cpu.rs b/sugarloaf/src/renderer/cpu.rs index e55e88c3..a4af5da2 100644 --- a/sugarloaf/src/renderer/cpu.rs +++ b/sugarloaf/src/renderer/cpu.rs @@ -377,13 +377,16 @@ pub fn render_cpu( }; buffer.fill(bg_u32); - // Grid pass: paint each panel's terminal cells (bg + glyphs) into - // the buffer before overlay vertices, so UI overlays composite on - // top. + // Grid passes: paint each panel's terminal cells (bg + glyphs) + // into the buffer before overlay vertices, so UI overlays + // composite on top. CPU backend doesn't handle kitty image layers + // (no image-overlay support on softbuffer), so the bg/text split + // here just runs back-to-back per panel. { let buf_slice: &mut [u32] = &mut buffer; for (grid, uniforms) in grids.iter() { - grid.render_cpu(buf_slice, ctx.width_px, ctx.height_px, uniforms); + grid.render_bg_cpu(buf_slice, ctx.width_px, ctx.height_px, uniforms); + grid.render_text_cpu(buf_slice, ctx.width_px, ctx.height_px, uniforms); } } diff --git a/sugarloaf/src/renderer/mod.rs b/sugarloaf/src/renderer/mod.rs index 9ebca6ba..b94206f5 100644 --- a/sugarloaf/src/renderer/mod.rs +++ b/sugarloaf/src/renderer/mod.rs @@ -758,15 +758,28 @@ pub struct ImageInstance { pub source_rect: [f32; 4], } -/// Which layer to render the image in (relative to text). +/// Which layer to render the image in. Mirrors ghostty's +/// three-bucket split (`renderer/image.zig:94-97`, +/// `renderer/generic.zig:1647-1695`): +/// +/// - `BelowBg` — `z < BG_LIMIT`. Drawn before the cell-bg pass; sits +/// underneath everything terminal-related. +/// - `BelowText` — `BG_LIMIT ≤ z < 0`. Drawn between cell-bg and +/// cell-text passes — the kitty default for "image with text on top". +/// - `AboveText` — `z >= 0`. Drawn after the cell-text pass; sits on +/// top of all glyphs. #[derive(Clone, Copy, PartialEq, Eq)] enum ImageLayer { - /// z < 0: rendered before the text pipeline. + BelowBg, BelowText, - /// z >= 0: rendered after the text pipeline. AboveText, } +/// Threshold separating `BelowBg` from `BelowText`. Matches ghostty's +/// `bg_limit = std.math.minInt(i32) / 2` at +/// `renderer/image.zig:377`. +const IMAGE_BG_LIMIT: i32 = i32::MIN / 2; + /// A single image draw command for the image pipeline. struct ImageDraw { image_id: u32, @@ -797,6 +810,17 @@ pub struct Renderer { /// Dedicated GPU texture for the background image, sized to the /// image dimensions instead of going through the glyph atlas. background_image_texture: Option, + /// Metal swap-chain state. One semaphore + one frame index for + /// the whole renderer regardless of how many split-pane grids + /// exist — mirrors ghostty's `SwapChain` at + /// `renderer/generic.zig:247`. Each render acquires one permit, + /// advances the index, hands the index to every grid's + /// `render_bg_metal` / `render_text_metal`, and releases the + /// permit from the command-buffer completion handler. + #[cfg(target_os = "macos")] + metal_frame_permits: crate::grid::metal::FramePermits, + #[cfg(target_os = "macos")] + metal_frame_index: usize, } /// Upload `pixels` to a fresh GPU texture using whatever backend `context` @@ -954,6 +978,10 @@ impl Renderer { image_draws: Vec::new(), background_image_dirty: None, background_image_texture: None, + #[cfg(target_os = "macos")] + metal_frame_permits: crate::grid::metal::new_frame_permits(), + #[cfg(target_os = "macos")] + metal_frame_index: 0, } } @@ -1432,7 +1460,9 @@ impl Renderer { dest_size: [overlay.width, overlay.height], source_rect: overlay.source_rect, }, - layer: if overlay.z_index < 0 { + layer: if overlay.z_index < IMAGE_BG_LIMIT { + ImageLayer::BelowBg + } else if overlay.z_index < 0 { ImageLayer::BelowText } else { ImageLayer::AboveText @@ -2170,6 +2200,17 @@ impl Renderer { } }; + // Acquire one swap-chain permit for the whole renderer. + // Blocks if 3 frames are already in flight — backpressure + // that keeps the CPU from outrunning the GPU. Mirrors + // ghostty's `SwapChain.nextFrame` at + // `renderer/generic.zig:295`. Single Arc, single permit, no + // matter how many split-pane grids the renderer is driving. + crate::grid::metal::acquire_frame_permit(&self.metal_frame_permits); + self.metal_frame_index = + (self.metal_frame_index + 1) % crate::grid::metal::FRAMES_IN_FLIGHT_PUB; + let frame = self.metal_frame_index; + loop { let instance_buffer = brush.instance_buffer_pool.lock().acquire(&context.device); @@ -2238,17 +2279,31 @@ impl Renderer { ) { return false; } - // Terminal grid passes — drawn after the window bg - // fill but before rich-text UI overlays, so each - // panel's cells composite over the window bg and - // under the tab-bar / assistant / search overlays. true })(); - if ok { - for (grid, uniforms) in grids.iter_mut() { - grid.render_metal(render_encoder, uniforms); - } - } + // Three-bucket image z-ordering — mirrors ghostty's + // `renderer/generic.zig:1640-1695`: + // + // bg fill / image (already drawn above) + // ↓ + // kitty z < BG_LIMIT (BelowBg) + // ↓ + // grid bg pass (per panel) + // ↓ + // kitty BG_LIMIT ≤ z < 0 (BelowText) + // ↓ + // grid text pass (per panel) + // ↓ + // kitty z >= 0 (AboveText) + // ↓ + // rich-text UI overlays (brush.render) + // ↓ + // UI text pass (labels) + // + // The two grid passes run inside a single iteration loop + // per panel — for multi-panel layouts the bg/text + // ordering is per-grid (one panel's text doesn't paint + // over another panel's bg image, and vice versa). let ok = ok && (|| { if has_images @@ -2257,7 +2312,7 @@ impl Renderer { &self.image_textures, brush, render_encoder, - ImageLayer::BelowText, + ImageLayer::BelowBg, &instance_buffer, &mut instance_offset, &globals, @@ -2265,18 +2320,26 @@ impl Renderer { { return false; } - if !brush.render( - &self.instances, - &self.vertices, - &self.draw_cmds, - &self.images, - render_encoder, - context, - &instance_buffer, - &mut instance_offset, - ) { + for (grid, uniforms) in grids.iter_mut() { + grid.render_bg_metal(render_encoder, frame, uniforms); + } + if has_images + && !Self::draw_images_metal( + &self.image_draws, + &self.image_textures, + brush, + render_encoder, + ImageLayer::BelowText, + &instance_buffer, + &mut instance_offset, + &globals, + ) + { return false; } + for (grid, uniforms) in grids.iter_mut() { + grid.render_text_metal(render_encoder, frame, uniforms); + } if has_images && !Self::draw_images_metal( &self.image_draws, @@ -2291,6 +2354,18 @@ impl Renderer { { return false; } + if !brush.render( + &self.instances, + &self.vertices, + &self.draw_cmds, + &self.images, + render_encoder, + context, + &instance_buffer, + &mut instance_offset, + ) { + return false; + } // UI text pass. Lazy-init on the first frame with // a Metal ctx; subsequent calls are no-ops. Runs // after brush.render / above-text images so UI @@ -2318,6 +2393,12 @@ impl Renderer { dropping frame", prev ); + // No completion handler will fire to release the + // swap-chain permit we acquired above — release + // it here so the next frame can run. + crate::grid::metal::release_frame_permit( + &self.metal_frame_permits, + ); return; } tracing::info!( @@ -2330,15 +2411,20 @@ impl Renderer { render_encoder.end_encoding(); - // Completion handler returns the buffer to the pool on GPU - // finish. The block fires on a Metal-internal thread; we - // hop into the pool's mutex to release. + // Completion handler returns the buffer to the pool + + // releases the swap-chain permit on GPU finish. The block + // fires on a Metal-internal thread; the `FramePermits` + // Arc inside the closure hops into its condvar to wake + // any frame waiting on `acquire_frame_permit`. One Arc + // clone per render — no per-grid duplication. let pool = brush.instance_buffer_pool.clone(); let buffer_cell = StdCell::new(Some(instance_buffer)); + let permits = self.metal_frame_permits.clone(); let block = ConcreteBlock::new(move |_cb: &metal::CommandBufferRef| { if let Some(b) = buffer_cell.take() { pool.lock().release(b); } + crate::grid::metal::release_frame_permit(&permits); }) .copy(); command_buffer.add_completed_handler(&block); diff --git a/sugarloaf/src/sugarloaf.rs b/sugarloaf/src/sugarloaf.rs index 4a796b60..0f6f6b09 100644 --- a/sugarloaf/src/sugarloaf.rs +++ b/sugarloaf/src/sugarloaf.rs @@ -1301,9 +1301,12 @@ impl Sugarloaf<'_> { } // Per-panel grid passes — draw cell backgrounds + grid text - // underneath everything else. + // underneath everything else. Vulkan doesn't yet interleave + // kitty image layers around the bg/text split — same as the + // wgpu path; follow-up. for (grid, uniforms) in grids.iter_mut() { - grid.render_vulkan(ctx, cmd, frame.slot, uniforms); + grid.render_bg_vulkan(ctx, cmd, frame.slot, uniforms); + grid.render_text_vulkan(ctx, cmd, frame.slot, uniforms); } // Rich-text quad pass — `Sugarloaf::quad()` / `rect()` calls @@ -1381,9 +1384,20 @@ impl Sugarloaf<'_> { }); // Grid passes first — cell bg/text composite under - // the rich-text UI overlays drawn below. + // the rich-text UI overlays drawn below. Wgpu + // doesn't yet interleave kitty image layers with + // the grid bg/text split (BrushRenderer::render + // owns kitty image draws inline), so for now the + // bg+text passes run back-to-back per panel — same + // visual result as the prior single render call + // and unchanged from ghostty's wgpu builds. + // Re-ordering kitty layers around the bg/text + // split would require pulling image draws out of + // BrushRenderer::render — Metal already does that; + // wgpu follow-up. for (grid, uniforms) in grids.iter_mut() { - grid.render_wgpu(&mut rpass, uniforms); + grid.render_bg_wgpu(&mut rpass, uniforms); + grid.render_text_wgpu(&mut rpass, uniforms); } self.renderer.render(ctx, &mut rpass); -- 2.51.2