diff --git a/.gitignore b/.gitignore index d2bc34ee..2b744d4c 100644 --- a/.gitignore +++ b/.gitignore @@ -26,3 +26,7 @@ misc/72/ # goreleaser dist/ + +# Compiled SPIR-V — generated at build time by sugarloaf/build.rs into +# OUT_DIR. The GLSL sources are checked in; the bytecode isn't. +sugarloaf/src/**/*.spv diff --git a/Cargo.lock b/Cargo.lock index 375aec4f..75768529 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -4769,6 +4769,7 @@ name = "sugarloaf" version = "0.4.0" dependencies = [ "approx", + "ash", "block", "bytemuck", "console_error_panic_hook", diff --git a/frontends/rioterm/src/application.rs b/frontends/rioterm/src/application.rs index eb3c901d..1b3c9066 100644 --- a/frontends/rioterm/src/application.rs +++ b/frontends/rioterm/src/application.rs @@ -306,7 +306,10 @@ impl ApplicationHandler for Application<'_> { } } } - RioEventType::Rio(RioEvent::UpdateGraphics { route_id: _, queues }) => { + RioEventType::Rio(RioEvent::UpdateGraphics { + route_id: _, + queues, + }) => { if let Some(route) = self.router.routes.get_mut(&window_id) { // Process graphics directly in sugarloaf let sugarloaf = &mut route.window.screen.sugarloaf; diff --git a/frontends/rioterm/src/screen/mod.rs b/frontends/rioterm/src/screen/mod.rs index 20b85b41..a3c85d41 100644 --- a/frontends/rioterm/src/screen/mod.rs +++ b/frontends/rioterm/src/screen/mod.rs @@ -35,7 +35,7 @@ use raw_window_handle::{RawDisplayHandle, RawWindowHandle}; use rio_backend::clipboard::Clipboard; use rio_backend::clipboard::ClipboardType; use rio_backend::config::layout::Margin; -use rio_backend::config::renderer::{Backend, Performance as RendererPerformance}; +use rio_backend::config::renderer::Backend; use rio_backend::crosswords::pos::{Boundary, CursorState, Direction, Line}; use rio_backend::crosswords::search::RegexSearch; use rio_backend::error::{RioError, RioErrorLevel, RioErrorType}; @@ -137,24 +137,40 @@ impl Screen<'_> { }, }; - let power_preference = match config.renderer.performance { - RendererPerformance::High => wgpu::PowerPreference::HighPerformance, - RendererPerformance::Low => wgpu::PowerPreference::LowPower, - }; - let backend = if config.renderer.use_cpu { SugarloafBackend::Cpu } else { match config.renderer.backend { Backend::Automatic => { - #[cfg(target_arch = "wasm32")] - let default_backend = - wgpu::Backends::BROWSER_WEBGPU | wgpu::Backends::GL; - #[cfg(not(target_arch = "wasm32"))] - let default_backend = wgpu::Backends::all(); - - SugarloafBackend::Wgpu(default_backend) + // Linux + macOS pick their native GPU backend (ash + // / Metal). Other targets fall back to the wgpu + // umbrella with whatever backends it can find. + #[cfg(target_os = "linux")] + { + SugarloafBackend::Vulkan + } + #[cfg(target_os = "macos")] + { + SugarloafBackend::Metal + } + #[cfg(not(any(target_os = "linux", target_os = "macos")))] + { + #[cfg(target_arch = "wasm32")] + let default_backend = + wgpu::Backends::BROWSER_WEBGPU | wgpu::Backends::GL; + #[cfg(not(target_arch = "wasm32"))] + let default_backend = wgpu::Backends::all(); + + SugarloafBackend::Wgpu(default_backend) + } } + // `Backend::Vulkan` from the user config now means the + // native ash backend on Linux. Other OSes fall through + // to the wgpu Vulkan path (Windows MSI, etc.) so the + // config option keeps working there. + #[cfg(target_os = "linux")] + Backend::Vulkan => SugarloafBackend::Vulkan, + #[cfg(not(target_os = "linux"))] Backend::Vulkan => SugarloafBackend::Wgpu(wgpu::Backends::VULKAN), Backend::GL => SugarloafBackend::Wgpu(wgpu::Backends::GL), Backend::WgpuMetal => SugarloafBackend::Wgpu(wgpu::Backends::METAL), @@ -165,7 +181,6 @@ impl Screen<'_> { }; let sugarloaf_renderer = SugarloafRenderer { - power_preference, backend, font_features: config.fonts.features.clone(), colorspace: config.window.colorspace.to_sugarloaf_colorspace(), diff --git a/rio-backend/src/config/mod.rs b/rio-backend/src/config/mod.rs index fcab4ff1..52b56842 100644 --- a/rio-backend/src/config/mod.rs +++ b/rio-backend/src/config/mod.rs @@ -596,9 +596,6 @@ impl Config { // Merge renderer fields individually if let Some(renderer_overwrite) = &platform_config.renderer { - if let Some(performance) = renderer_overwrite.performance { - self.renderer.performance = performance; - } if let Some(backend) = &renderer_overwrite.backend { self.renderer.backend = backend.clone(); } diff --git a/rio-backend/src/config/platform.rs b/rio-backend/src/config/platform.rs index cb105f0e..54c0e827 100644 --- a/rio-backend/src/config/platform.rs +++ b/rio-backend/src/config/platform.rs @@ -112,8 +112,6 @@ pub struct PlatformNavigation { /// Platform-specific renderer config with optional fields for selective override #[derive(Default, Debug, Serialize, Deserialize, PartialEq, Clone)] pub struct PlatformRenderer { - #[serde(default = "Option::default")] - pub performance: Option, #[serde(default = "Option::default")] pub backend: Option, #[serde(default = "Option::default", rename = "disable-unfocused-render")] diff --git a/rio-backend/src/config/renderer.rs b/rio-backend/src/config/renderer.rs index cfc8fc26..004ddd9c 100644 --- a/rio-backend/src/config/renderer.rs +++ b/rio-backend/src/config/renderer.rs @@ -4,8 +4,6 @@ use sugarloaf::Filter; #[derive(Debug, Clone, PartialEq, Deserialize, Serialize)] pub struct Renderer { - #[serde(default = "Performance::default")] - pub performance: Performance, #[serde(default = "Backend::default", skip_serializing)] pub backend: Backend, #[serde(default = "bool::default", rename = "disable-unfocused-render")] @@ -60,7 +58,6 @@ impl RendererStategy { impl Default for Renderer { fn default() -> Renderer { Renderer { - performance: Performance::default(), backend: Backend::default(), disable_unfocused_render: false, disable_occluded_render: default_disable_occluded_render(), @@ -71,28 +68,6 @@ impl Default for Renderer { } } -#[derive(Default, Debug, Serialize, Deserialize, PartialEq, Clone, Copy)] -pub enum Performance { - #[default] - #[serde(alias = "high")] - High, - #[serde(alias = "low")] - Low, -} - -impl Display for Performance { - fn fmt(&self, f: &mut std::fmt::Formatter) -> std::fmt::Result { - match self { - Performance::High => { - write!(f, "High") - } - Performance::Low => { - write!(f, "Low") - } - } - } -} - #[derive(Debug, Default, Serialize, Deserialize, Clone, PartialEq)] pub enum Backend { // Leave Sugarloaf/WGPU to decide diff --git a/rio-window/src/platform_impl/input_rate.rs b/rio-window/src/platform_impl/input_rate.rs index e96776f0..f59bc6df 100644 --- a/rio-window/src/platform_impl/input_rate.rs +++ b/rio-window/src/platform_impl/input_rate.rs @@ -23,7 +23,7 @@ use std::time::{Duration, Instant}; #[derive(Debug, Clone)] -pub(crate) struct InputRateTracker { +pub struct InputRateTracker { timestamps: Vec, window: Duration, inputs_per_second: u32, @@ -60,8 +60,7 @@ impl InputRateTracker { // `min_events` = inputs_per_second × window_ms / 1000. For // defaults (60/s, 100 ms) that's 6 events in the last 100 ms. - let min_events = - self.inputs_per_second as u128 * self.window.as_millis() / 1000; + let min_events = self.inputs_per_second as u128 * self.window.as_millis() / 1000; if self.timestamps.len() as u128 >= min_events { self.sustain_until = now + self.sustain_duration; } diff --git a/rio-window/src/platform_impl/linux/mod.rs b/rio-window/src/platform_impl/linux/mod.rs index 56a9d63b..88c66770 100644 --- a/rio-window/src/platform_impl/linux/mod.rs +++ b/rio-window/src/platform_impl/linux/mod.rs @@ -854,6 +854,12 @@ impl EventLoopProxy { } } +// X11's ActiveEventLoop is much larger than the Wayland variant; the +// imbalance is intentional (X11 carries more state) and only one +// variant is constructed per process. Boxing the larger variant +// would just trade stack size for a heap allocation that lives the +// whole event-loop lifetime. +#[allow(clippy::large_enum_variant)] pub enum ActiveEventLoop { #[cfg(wayland_platform)] Wayland(Box), diff --git a/rio-window/src/platform_impl/linux/wayland/state.rs b/rio-window/src/platform_impl/linux/wayland/state.rs index 524b50c8..843d3b49 100644 --- a/rio-window/src/platform_impl/linux/wayland/state.rs +++ b/rio-window/src/platform_impl/linux/wayland/state.rs @@ -120,9 +120,8 @@ pub struct WinitState { /// input (≥ 60 events/sec over 100 ms) extends the 1-second /// presentation window — single keystrokes don't. See /// `platform_impl::input_rate`. - pub input_rate_tracker: std::cell::RefCell< - crate::platform_impl::input_rate::InputRateTracker, - >, + pub input_rate_tracker: + std::cell::RefCell, } impl WinitState { diff --git a/rio-window/src/platform_impl/linux/x11/mod.rs b/rio-window/src/platform_impl/linux/x11/mod.rs index 80b023f5..6fb12151 100644 --- a/rio-window/src/platform_impl/linux/x11/mod.rs +++ b/rio-window/src/platform_impl/linux/x11/mod.rs @@ -149,8 +149,7 @@ pub struct ActiveEventLoop { /// input (≥ 60 events/sec over 100 ms) extends the 1-second /// presentation window — single keystrokes don't. See /// `platform_impl::input_rate`. - input_rate_tracker: - RefCell, + input_rate_tracker: RefCell, /// Shared with `EventLoopState` and every window. Set by /// `request_redraw`, checked by `has_pending`, cleared by /// the vsync tick after fanning out. diff --git a/sugarloaf/Cargo.toml b/sugarloaf/Cargo.toml index 0aef6605..e2b44a98 100644 --- a/sugarloaf/Cargo.toml +++ b/sugarloaf/Cargo.toml @@ -4,8 +4,10 @@ version.workspace = true edition.workspace = true license.workspace = true authors = ["Raphael Amorim "] +build = "build.rs" include = [ "Cargo.toml", + "build.rs", "src/**/*.ttf", "src/**/*.otf", "src/**/*.wgsl", @@ -15,6 +17,7 @@ include = [ "src/components/filters/**/*.png", "src/**/*.rs", "src/**/*.metal", + "src/**/*.glsl", ] description = "Sugarloaf is Rio rendering engine, designed to be multiplatform. It is based on WebGPU, Rust library for Desktops and WebAssembly for Web (JavaScript). This project is created and maintained for Rio terminal purposes but feel free to use it." documentation = "https://docs.rs/crate/sugarloaf/latest" @@ -34,7 +37,12 @@ targets = [ ] [dependencies] -wgpu = { workspace = true } +# `wgpu` and the `librashader-*` filter chain are gated behind the +# `wgpu` feature. Default-off on macOS + Linux (the native Metal / +# Vulkan backends are preferred and skip the wgpu translation +# layer). Windows + WASM enable it via target-specific deps in the +# downstream crate (`frontends/rioterm/Cargo.toml`). +wgpu = { workspace = true, optional = true } bytemuck = { workspace = true } tracing = { workspace = true } serde = { workspace = true, features = ["derive"] } @@ -59,14 +67,15 @@ futures = { workspace = true } tiny-skia = "0.12.0" wide = "1.2.0" -librashader = { version = "0.10.0", default-features = false, features = ["runtime-wgpu", "stable", "presets"] } -librashader-common = "0.10" -librashader-presets = "0.10" -librashader-preprocess = "0.10" -librashader-pack = "0.10" -librashader-reflect = { version = "0.10", features = ["stable", "wgsl"], default-features = false } -librashader-runtime = "0.10" -librashader-cache = "0.10" +# librashader is wgpu-only — gate the entire filter chain together. +librashader = { version = "0.10.0", default-features = false, features = ["runtime-wgpu", "stable", "presets"], optional = true } +librashader-common = { version = "0.10", optional = true } +librashader-presets = { version = "0.10", optional = true } +librashader-preprocess = { version = "0.10", optional = true } +librashader-pack = { version = "0.10", optional = true } +librashader-reflect = { version = "0.10", features = ["stable", "wgsl"], default-features = false, optional = true } +librashader-runtime = { version = "0.10", optional = true } +librashader-cache = { version = "0.10", optional = true } thiserror = "2.0.1" [target.'cfg(target_os = "macos")'.dependencies] @@ -93,6 +102,11 @@ ttf-parser = { version = "0.25.1", default-features = false, features = ["std", [target.'cfg(all(unix, not(any(target_os = "macos", target_os = "android"))))'.dependencies] fontconfig-parser = { version = "0.5.8", default-features = false } +# Native Vulkan backend (mirrors the Metal backend in scope: no librashader +# filters, owns its own swapchain + pipelines). Linux-only for now — Windows +# would need khr::win32_surface added to context::vulkan, and macOS keeps +# the native Metal path. +ash = "0.38" [dev-dependencies] rio-window = { workspace = true } @@ -104,6 +118,21 @@ criterion = { workspace = true } default = ["scale", "render"] scale = ["yazi"] render = ["scale", "zeno/eval"] +# Pull in `wgpu` + the librashader filter chain. Required on Windows +# and WASM (no native backend yet); optional on Linux + macOS where +# the native Vulkan / Metal backends cover everything except the +# librashader CRT/scanline filters. +wgpu = [ + "dep:wgpu", + "dep:librashader", + "dep:librashader-common", + "dep:librashader-presets", + "dep:librashader-preprocess", + "dep:librashader-pack", + "dep:librashader-reflect", + "dep:librashader-runtime", + "dep:librashader-cache", +] [target.'cfg(target_arch = "wasm32")'.dependencies] console_error_panic_hook = "0.1.7" diff --git a/sugarloaf/README.md b/sugarloaf/README.md index e2eab5db..4dd72d0a 100644 --- a/sugarloaf/README.md +++ b/sugarloaf/README.md @@ -6,6 +6,22 @@ Sugarloaf is Rio rendering engine, designed to be multiplatform. It is based on cargo run --example text ``` +## Build dependencies + +### Linux — Vulkan backend + +The native Vulkan backend (default on Linux) compiles its GLSL shaders to SPIR-V at build time. You need one GLSL → SPIR-V compiler installed on the build host: + +| Distro | Command | +|---|---| +| Debian / Ubuntu | `apt install glslang-tools` (or `apt install glslc`) | +| Arch | `pacman -S shaderc` (provides `glslc`) | +| Fedora | `dnf install glslang` (or `dnf install glslc`) | + +`glslc` is preferred when both are present. Override with `GLSLC=/path/to/binary` or `GLSLANG_VALIDATOR=/path/to/binary`. + +The compiled SPIR-V lives in `OUT_DIR` per build — the source `.glsl` files are checked in but the `.spv` artifacts are gitignored. + ## WASM Tests ### Setup diff --git a/sugarloaf/build.rs b/sugarloaf/build.rs new file mode 100644 index 00000000..f3b2cb0e --- /dev/null +++ b/sugarloaf/build.rs @@ -0,0 +1,257 @@ +// Copyright (c) 2023-present, Raphael Amorim. +// +// This source code is licensed under the MIT license found in the +// LICENSE file in the root directory of this source tree. + +//! Compile sugarloaf's GLSL shaders to SPIR-V at build time. +//! +//! The Vulkan backend is Linux-only — on other targets this script is +//! a no-op. On Linux we walk a hard-coded list of `.glsl` files +//! (kept in lock-step with the `include_bytes!` call sites) and shell +//! out to `glslc` (preferred) or `glslangValidator` (Debian fallback) +//! to compile each into `$OUT_DIR/.spv`. Compiled bytes get +//! pulled in via `include_bytes!(concat!(env!("OUT_DIR"), "/..."))` +//! at the call sites, so the `.spv` files never live in the source +//! tree (and never need to be committed). +//! +//! ## Required tooling +//! +//! - **Debian**: `apt install glslang-tools` (provides +//! `glslangValidator`) or `apt install glslc` (provides Google's +//! `glslc`, available since Bookworm). +//! - **Arch**: `pacman -S shaderc` (provides `glslc`). +//! - **macOS / Windows**: not required — the Vulkan backend isn't +//! built on those targets, so this script returns early. +//! +//! Override the compiler with `GLSLC=/path/to/binary` if your +//! compiler isn't on PATH or you want a specific version. +//! +//! ## TODO: drop the system-binary requirement +//! +//! `librust-naga-dev` (the wgpu shader translator, pure Rust) has a +//! GLSL frontend + SPIR-V backend and would let us compile shaders +//! in-process with zero system tooling. As of 2026-04 it's only in +//! Debian forky/sid (24.0.0-3); not yet in trixie/stable. Once it +//! lands in Debian stable, switch to a `naga` build-dependency and +//! delete the `glslc`/`glslangValidator` subprocess plumbing. + +use std::path::{Path, PathBuf}; +use std::process::Command; + +/// Directories we scan for `*.{vert,frag}.glsl` sources. New shader +/// files dropped into either are picked up automatically — no +/// build.rs edit needed. File stems must be unique across all +/// scanned directories (we flatten output into a single `OUT_DIR`); +/// `cargo build` panics with a clear message if two sources end up +/// with the same `.spv` name. +/// +/// Same auto-discovery pattern as `adrien-ben/vulkan-tutorial-rs`'s +/// `build.rs` — a `fs::read_dir` walk, std-only, no `walkdir` / +/// `glob` dep. Survey of comparable Rust+ash projects (kajiya, +/// screen-13, vulkano-shaders, etc.) showed hard-coded lists are the +/// norm; this is the cleanest "drop a file and it works" alternative +/// that doesn't pull in a new crate. +const SHADER_DIRS: &[&str] = &[ + // renderer (rich-text quad / non-quad / image / bootstrap) + "src/renderer", + // grid (per-panel terminal cell + text + UI text overlay) + "src/grid/shaders", +]; + +fn main() { + let target_os = std::env::var("CARGO_CFG_TARGET_OS").unwrap_or_default(); + if target_os != "linux" { + // Vulkan backend is Linux-only; nothing to compile. We still + // emit `rerun-if-env-changed` so a future port (e.g. + // `khr::win32_surface`) flipping target conditions takes + // effect on the next cargo invocation. + println!("cargo:rerun-if-env-changed=CARGO_CFG_TARGET_OS"); + return; + } + + println!("cargo:rerun-if-env-changed=GLSLC"); + println!("cargo:rerun-if-env-changed=GLSLANG_VALIDATOR"); + println!("cargo:rerun-if-changed=build.rs"); + + let out_dir = + PathBuf::from(std::env::var_os("OUT_DIR").expect("OUT_DIR must be set by cargo")); + + let compiler = locate_compiler(); + eprintln!("sugarloaf build.rs: GLSL compiler = {:?}", compiler); + + let sources = discover_shaders(SHADER_DIRS); + if sources.is_empty() { + // Survives a partial crate checkout (no source dir present) + // without breaking the build — but warn so the missing dir + // is visible in cargo output. + println!( + "cargo:warning=sugarloaf build.rs: no GLSL shaders found in {SHADER_DIRS:?}" + ); + return; + } + + for src in &sources { + println!("cargo:rerun-if-changed={}", src.display()); + compile(&compiler, src, &out_dir); + } +} + +/// Walk `dirs` looking for `*.vert.glsl` and `*.frag.glsl` files. +/// Returns absolute-relative-to-CARGO_MANIFEST_DIR paths. Also +/// emits `cargo:rerun-if-changed` for each *directory* so cargo +/// triggers a rebuild when a new file is added (without that, a new +/// `.glsl` would only be picked up after a `cargo clean`). +fn discover_shaders(dirs: &[&str]) -> Vec { + use std::collections::HashSet; + + let mut out: Vec = Vec::new(); + let mut seen_stems: HashSet = HashSet::new(); + + for dir in dirs { + // Tell cargo to rebuild when files appear/disappear in this + // directory. (cargo also implicitly tracks individual files + // we list via rerun-if-changed below; this catches NEW files.) + println!("cargo:rerun-if-changed={dir}"); + + let entries = match std::fs::read_dir(dir) { + Ok(e) => e, + Err(_) => continue, + }; + for entry in entries.flatten() { + let path = entry.path(); + let name = match path.file_name().and_then(|s| s.to_str()) { + Some(n) => n.to_owned(), + None => continue, + }; + if !(name.ends_with(".vert.glsl") || name.ends_with(".frag.glsl")) { + continue; + } + let stem = name + .strip_suffix(".glsl") + .expect("just checked the suffix") + .to_owned(); + if !seen_stems.insert(stem.clone()) { + panic!( + "sugarloaf build.rs: duplicate shader stem `{stem}` across \ + scanned dirs — `OUT_DIR` would collide. Rename one of them." + ); + } + out.push(path); + } + } + + // Stable order so build logs and rerun-if-changed lines diff + // cleanly across runs. + out.sort(); + out +} + +#[derive(Debug)] +enum Compiler { + /// Google's `glslc` from shaderc — better diagnostics, what + /// almost every Vulkan tutorial uses. + Glslc(PathBuf), + /// Khronos reference compiler. The `-V` (Vulkan) flag is + /// load-bearing: without it we'd get GL SPIR-V which Vulkan + /// drivers reject. + Glslang(PathBuf), +} + +/// PATH-like lookup with `GLSLC` / `GLSLANG_VALIDATOR` env-var +/// override. We don't pull in the `which` crate (Debian doesn't +/// universally package it) — the lookup is a few lines. +fn locate_compiler() -> Compiler { + if let Some(p) = env_path("GLSLC") { + return Compiler::Glslc(p); + } + if let Some(p) = env_path("GLSLANG_VALIDATOR") { + return Compiler::Glslang(p); + } + if let Some(p) = which("glslc") { + return Compiler::Glslc(p); + } + if let Some(p) = which("glslangValidator") { + return Compiler::Glslang(p); + } + panic!( + "\nsugarloaf: no GLSL → SPIR-V compiler found.\n\ + Install one of:\n \ + * Debian: `apt install glslang-tools` (provides glslangValidator)\n \ + * or `apt install glslc` (Google's glslc, Bookworm+)\n \ + * Arch: `pacman -S shaderc` (provides glslc)\n \ + * Or set GLSLC=/path/to/binary or GLSLANG_VALIDATOR=/path/to/binary.\n" + ); +} + +fn env_path(name: &str) -> Option { + let v = std::env::var_os(name)?; + if v.is_empty() { + return None; + } + let p = PathBuf::from(v); + if p.is_file() { + Some(p) + } else { + None + } +} + +fn which(binary: &str) -> Option { + let path = std::env::var_os("PATH")?; + for dir in std::env::split_paths(&path) { + let candidate = dir.join(binary); + if candidate.is_file() { + return Some(candidate); + } + } + None +} + +fn compile(compiler: &Compiler, src: &Path, out_dir: &Path) { + let src_str = src.to_string_lossy(); + let stage = if src_str.ends_with(".vert.glsl") { + "vertex" + } else if src_str.ends_with(".frag.glsl") { + "fragment" + } else { + panic!("unrecognised GLSL stage suffix in {src_str}"); + }; + + let stem = src + .file_name() + .expect("source path has a file name") + .to_string_lossy(); + // foo.vert.glsl → foo.vert.spv + let spv_name = stem + .strip_suffix(".glsl") + .expect("source ends in .glsl") + .to_string() + + ".spv"; + let dst = out_dir.join(&spv_name); + + let output = match compiler { + Compiler::Glslc(bin) => Command::new(bin) + .arg(format!("-fshader-stage={stage}")) + .arg(src) + .arg("-o") + .arg(&dst) + .output(), + Compiler::Glslang(bin) => Command::new(bin) + .arg("-V") + .arg("-S") + .arg(stage) + .arg(src) + .arg("-o") + .arg(&dst) + .output(), + } + .unwrap_or_else(|e| panic!("failed to invoke {compiler:?} on {src_str}: {e}")); + + if !output.status.success() { + let stdout = String::from_utf8_lossy(&output.stdout); + let stderr = String::from_utf8_lossy(&output.stderr); + panic!( + "GLSL compile failed for {src_str}:\n--stdout--\n{stdout}\n--stderr--\n{stderr}" + ); + } +} diff --git a/sugarloaf/src/components/core/mod.rs b/sugarloaf/src/components/core/mod.rs index f84865a8..2cceaf63 100644 --- a/sugarloaf/src/components/core/mod.rs +++ b/sugarloaf/src/components/core/mod.rs @@ -1,6 +1,9 @@ pub mod image; pub mod uniforms; // pub mod svg; +// `buffer` is a thin wrapper over `wgpu::Buffer` — gated together +// with the rest of the wgpu code. +#[cfg(feature = "wgpu")] pub mod buffer; pub mod shapes; diff --git a/sugarloaf/src/components/mod.rs b/sugarloaf/src/components/mod.rs index a0f95bdf..22336a9f 100644 --- a/sugarloaf/src/components/mod.rs +++ b/sugarloaf/src/components/mod.rs @@ -1,2 +1,7 @@ pub mod core; +// `filters` is the librashader integration — wgpu-only by upstream +// design (`librashader-runtime-wgpu`). Gated together with the rest +// of the wgpu code so the dep tree drops cleanly on Linux/macOS +// builds that don't enable the `wgpu` feature. +#[cfg(feature = "wgpu")] pub mod filters; diff --git a/sugarloaf/src/context/mod.rs b/sugarloaf/src/context/mod.rs index b7b8ce65..408f03de 100644 --- a/sugarloaf/src/context/mod.rs +++ b/sugarloaf/src/context/mod.rs @@ -1,6 +1,9 @@ pub mod cpu; #[cfg(target_os = "macos")] pub mod metal; +#[cfg(target_os = "linux")] +pub mod vulkan; +#[cfg(feature = "wgpu")] pub mod webgpu; use crate::sugarloaf::{SugarloafBackend, SugarloafWindow}; @@ -12,10 +15,20 @@ pub struct Context<'a> { #[allow(clippy::large_enum_variant)] pub enum ContextType<'a> { + #[cfg(feature = "wgpu")] Wgpu(webgpu::WgpuContext<'a>), #[cfg(target_os = "macos")] Metal(metal::MetalContext), + #[cfg(target_os = "linux")] + Vulkan(vulkan::VulkanContext), Cpu(cpu::CpuContext), + /// Lifetime placeholder for the Wgpu variant when it's + /// feature-gated out — keeps `'a` referenced across the enum so + /// the compiler doesn't complain about an unused parameter on + /// builds without wgpu. + #[cfg(not(feature = "wgpu"))] + #[doc(hidden)] + _Phantom(std::marker::PhantomData<&'a ()>), } impl Context<'_> { @@ -24,6 +37,7 @@ impl Context<'_> { renderer_config: SugarloafRenderer, ) -> Context<'a> { let inner = match renderer_config.backend { + #[cfg(feature = "wgpu")] SugarloafBackend::Wgpu(backends) => ContextType::Wgpu( webgpu::WgpuContext::new(sugarloaf_window, renderer_config, backends), ), @@ -31,6 +45,10 @@ impl Context<'_> { SugarloafBackend::Metal => { ContextType::Metal(metal::MetalContext::new(sugarloaf_window)) } + #[cfg(target_os = "linux")] + SugarloafBackend::Vulkan => { + ContextType::Vulkan(vulkan::VulkanContext::new(sugarloaf_window)) + } SugarloafBackend::Cpu => { ContextType::Cpu(cpu::CpuContext::new(sugarloaf_window)) } @@ -42,16 +60,22 @@ impl Context<'_> { #[inline] pub fn scale(&self) -> f32 { match &self.inner { + #[cfg(feature = "wgpu")] ContextType::Wgpu(ctx) => ctx.scale, #[cfg(target_os = "macos")] ContextType::Metal(ctx) => ctx.scale, + #[cfg(target_os = "linux")] + ContextType::Vulkan(ctx) => ctx.scale, ContextType::Cpu(ctx) => ctx.scale, + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), } } #[inline] pub fn set_scale(&mut self, scale: f32) { match &mut self.inner { + #[cfg(feature = "wgpu")] ContextType::Wgpu(ctx) => { ctx.set_scale(scale); } @@ -59,19 +83,30 @@ impl Context<'_> { ContextType::Metal(ctx) => { ctx.set_scale(scale); } + #[cfg(target_os = "linux")] + ContextType::Vulkan(ctx) => { + ctx.set_scale(scale); + } ContextType::Cpu(ctx) => { ctx.set_scale(scale); } + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), } } #[inline] pub fn size(&self) -> SugarloafWindowSize { match &self.inner { + #[cfg(feature = "wgpu")] ContextType::Wgpu(ctx) => ctx.size, #[cfg(target_os = "macos")] ContextType::Metal(ctx) => ctx.size, + #[cfg(target_os = "linux")] + ContextType::Vulkan(ctx) => ctx.size, ContextType::Cpu(ctx) => ctx.size, + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), } } @@ -81,20 +116,30 @@ impl Context<'_> { } match &mut self.inner { + #[cfg(feature = "wgpu")] ContextType::Wgpu(ctx) => ctx.resize(width, height), #[cfg(target_os = "macos")] ContextType::Metal(ctx) => ctx.resize(width, height), + #[cfg(target_os = "linux")] + ContextType::Vulkan(ctx) => ctx.resize(width, height), ContextType::Cpu(ctx) => ctx.resize(width, height), + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), } } #[inline] pub fn supports_f16(&self) -> bool { match &self.inner { + #[cfg(feature = "wgpu")] ContextType::Wgpu(ctx) => ctx.supports_f16(), #[cfg(target_os = "macos")] ContextType::Metal(ctx) => ctx.supports_f16(), + #[cfg(target_os = "linux")] + ContextType::Vulkan(ctx) => ctx.supports_f16(), ContextType::Cpu(ctx) => ctx.supports_f16(), + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), } } } diff --git a/sugarloaf/src/context/vulkan.rs b/sugarloaf/src/context/vulkan.rs new file mode 100644 index 00000000..5f671357 --- /dev/null +++ b/sugarloaf/src/context/vulkan.rs @@ -0,0 +1,1441 @@ +//! Vulkan backend built directly on `ash`. +//! +//! Mirrors `context::metal::MetalContext` in shape and intent: one struct +//! owning everything needed to present a swapchain image, plus the +//! per-frame synchronisation primitives. No wgpu involvement. +//! +//! Targets Vulkan 1.3 so we can reach for `VK_KHR_dynamic_rendering` (core +//! in 1.3) later without changing device creation. Anyone on a driver +//! older than early-2022 will fail at `create_instance` with +//! `ERROR_INCOMPATIBLE_DRIVER` — same class of failure as an ancient GPU +//! on the Metal path. +//! +//! Surface creation dispatches on `raw-window-handle` variants inline +//! rather than pulling in `ash-window` — that crate is not in Debian and +//! `ash-window` buys us ~30 lines of glue per platform that we'd rather +//! own. + +use crate::sugarloaf::{Colorspace, SugarloafWindow, SugarloafWindowSize}; +use ash::khr; +use ash::vk; +use ash::{Device, Entry, Instance}; +use raw_window_handle::{ + HasDisplayHandle, HasWindowHandle, RawDisplayHandle, RawWindowHandle, +}; +use std::ffi::{c_char, CStr}; + +/// How many frames the CPU is allowed to pipeline ahead of the GPU. +/// Three matches the Metal backend (`MetalLayer::set_maximum_drawable_count(3)` +/// in `context::metal::MetalContext::new`) and Apple's standard sample +/// pattern — CPU / GPU / compositor each work on their own frame in +/// parallel. The cost is two extra `FrameSync` slots and one extra +/// swapchain image's worth of memory. +pub const FRAMES_IN_FLIGHT: usize = 3; + +/// One set of synchronisation objects + a command pool & pre-allocated +/// primary buffer, reused each time the same slot comes around. The +/// `in_flight` fence is signalled by the submit that uses this slot so +/// the *next* owner knows the GPU is done with this slot's resources. +struct FrameSync { + image_available: vk::Semaphore, + render_finished: vk::Semaphore, + in_flight: vk::Fence, + cmd_pool: vk::CommandPool, + cmd_buffer: vk::CommandBuffer, +} + +pub struct VulkanContext { + // Logical fields for the public surface. + pub size: SugarloafWindowSize, + pub scale: f32, + pub supports_f16: bool, + pub colorspace: Colorspace, + /// Updated on every `acquire_frame()` — `true` if the driver hinted + /// that the swapchain is out of date and we should recreate at our + /// earliest convenience (we already did if ERROR_OUT_OF_DATE_KHR, but + /// SUBOPTIMAL_KHR says "still usable this frame"). + pub needs_recreate: bool, + + // Per-frame state. + frame_index: usize, + frames: [FrameSync; FRAMES_IN_FLIGHT], + + // Swapchain state. Rebuilt by `resize()`. + swapchain_extent: vk::Extent2D, + swapchain_color_space: vk::ColorSpaceKHR, + swapchain_format: vk::Format, + swapchain_images: Vec, + swapchain_views: Vec, + swapchain: vk::SwapchainKHR, + swapchain_loader: khr::swapchain::Device, + + // Core device. + queue: vk::Queue, + // Kept around so future phases (atlas uploads, a dedicated transfer + // pool, pipeline creation) don't have to re-probe the family. + #[allow(dead_code)] + queue_family_index: u32, + device: Device, + physical_device: vk::PhysicalDevice, + + /// Pipeline cache shared by every `create_graphics_pipelines` + /// call. Loaded from `~/.cache/rio/sugarloaf-vulkan.cache` (best + /// effort) at startup and serialised back on `Drop`. Saves + /// ~10–50ms of pipeline build time on subsequent launches. + pipeline_cache: vk::PipelineCache, + + // Instance-level state — held last so it outlives everything above in + // the Drop impl (drop order = declaration order). + surface: vk::SurfaceKHR, + surface_loader: khr::surface::Instance, + /// Debug-utils messenger, present only when validation layers + /// were requested via `RIO_VULKAN_VALIDATION=1`. Drops before + /// `instance` (declaration order) so the messenger handle is + /// destroyed while the instance is still alive. + _debug_messenger: Option, + instance: Instance, + _entry: Entry, +} + +/// Owns one `vk::DebugUtilsMessengerEXT` and its loader. Destroyed +/// in `Drop` — the loader needs the parent `Instance` to still be +/// valid, which the field-order convention ensures. +struct DebugMessenger { + loader: ash::ext::debug_utils::Instance, + handle: vk::DebugUtilsMessengerEXT, +} + +impl Drop for DebugMessenger { + fn drop(&mut self) { + unsafe { + self.loader.destroy_debug_utils_messenger(self.handle, None); + } + } +} + +/// In-flight handle returned by `acquire_frame()`. The caller records +/// commands into `cmd_buffer` targeting `image` / `image_view`, then +/// hands it back to `present_frame()`. +pub struct VulkanFrame { + pub image_index: u32, + pub image: vk::Image, + pub image_view: vk::ImageView, + pub cmd_buffer: vk::CommandBuffer, + pub extent: vk::Extent2D, + pub format: vk::Format, + /// Frame-in-flight slot for this frame. Renderers (grid, text, + /// images) ring their per-frame GPU resources by this index; the + /// `in_flight` fence wait inside `acquire_frame` proved this slot + /// is GPU-idle, so writing into slot `N`'s buffers from the CPU + /// is safe. + pub slot: usize, +} + +impl VulkanContext { + pub fn new(sugarloaf_window: SugarloafWindow) -> Self { + let size = sugarloaf_window.size; + let scale = sugarloaf_window.scale; + + // Loading the loader itself can fail if libvulkan.so is missing — + // which is the expected failure on a machine without a Vulkan + // driver installed. We let the panic propagate: the caller is + // `Context::new` and the backend selection happened upstream, so + // there's no graceful degradation path here (the WGPU/CPU + // backends live behind different enum variants). + let entry = + unsafe { Entry::load() }.expect("failed to load Vulkan loader (libvulkan)"); + + let validation_requested = validation_requested(); + let instance = create_instance(&entry, &sugarloaf_window, validation_requested); + let _debug_messenger = if validation_requested { + create_debug_messenger(&entry, &instance) + } else { + None + }; + let surface_loader = khr::surface::Instance::new(&entry, &instance); + let surface = create_surface(&entry, &instance, &sugarloaf_window); + + let (physical_device, queue_family_index) = + pick_physical_device(&instance, &surface_loader, surface); + + let device = create_device(&instance, physical_device, queue_family_index); + let queue = unsafe { device.get_device_queue(queue_family_index, 0) }; + let pipeline_cache = create_pipeline_cache(&device); + + let swapchain_loader = khr::swapchain::Device::new(&instance, &device); + + let ( + swapchain, + swapchain_format, + swapchain_color_space, + swapchain_extent, + swapchain_images, + swapchain_views, + ) = create_swapchain( + &device, + &surface_loader, + &swapchain_loader, + physical_device, + surface, + size.width as u32, + size.height as u32, + vk::SwapchainKHR::null(), + ); + + let frames = create_frames(&device, queue_family_index); + + // f16 = Vulkan's VK_KHR_shader_float16_int8 feature. Probe at + // device creation time in a follow-up; for the MVP we report + // false, matching the conservative default. + let supports_f16 = false; + + tracing::info!( + "Vulkan device created: {}", + physical_device_name(&instance, physical_device) + ); + tracing::info!( + "Swapchain: {:?} {}x{} ({} images)", + swapchain_format, + swapchain_extent.width, + swapchain_extent.height, + swapchain_images.len() + ); + log_memory_heap_choice(&instance, physical_device); + + VulkanContext { + size, + scale, + supports_f16, + colorspace: Colorspace::Srgb, + needs_recreate: false, + frame_index: 0, + frames, + swapchain_extent, + swapchain_color_space, + swapchain_format, + swapchain_images, + swapchain_views, + swapchain, + swapchain_loader, + queue, + queue_family_index, + device, + physical_device, + pipeline_cache, + surface, + surface_loader, + _debug_messenger, + instance, + _entry: entry, + } + } + + #[inline] + pub fn set_scale(&mut self, scale: f32) { + self.scale = scale; + } + + #[inline] + pub fn supports_f16(&self) -> bool { + self.supports_f16 + } + + pub fn resize(&mut self, width: u32, height: u32) { + if width == 0 || height == 0 { + return; + } + self.size.width = width as f32; + self.size.height = height as f32; + self.recreate_swapchain(width, height); + } + + fn recreate_swapchain(&mut self, width: u32, height: u32) { + // The spec requires no resources tied to the old swapchain be + // in use. Easiest safe path: wait for the device to go idle. + // This is a resize, not a per-frame operation, so the stall is + // acceptable (wgpu does the same thing). + unsafe { + let _ = self.device.device_wait_idle(); + } + + for &view in &self.swapchain_views { + unsafe { self.device.destroy_image_view(view, None) }; + } + self.swapchain_views.clear(); + self.swapchain_images.clear(); + + let old = self.swapchain; + let (swapchain, format, color_space, extent, images, views) = create_swapchain( + &self.device, + &self.surface_loader, + &self.swapchain_loader, + self.physical_device, + self.surface, + width, + height, + old, + ); + + unsafe { self.swapchain_loader.destroy_swapchain(old, None) }; + + self.swapchain = swapchain; + self.swapchain_format = format; + self.swapchain_color_space = color_space; + self.swapchain_extent = extent; + self.swapchain_images = images; + self.swapchain_views = views; + self.needs_recreate = false; + } + + /// Acquire the next swapchain image and begin the per-frame command + /// buffer. Returns `None` if the swapchain needed recreation (caller + /// should skip this frame). + pub fn acquire_frame(&mut self) -> Option { + let slot = self.frame_index; + let sync = &self.frames[slot]; + + unsafe { + self.device + .wait_for_fences(&[sync.in_flight], true, u64::MAX) + .expect("wait_for_fences"); + } + + let (image_index, suboptimal) = unsafe { + match self.swapchain_loader.acquire_next_image( + self.swapchain, + u64::MAX, + sync.image_available, + vk::Fence::null(), + ) { + Ok(pair) => pair, + Err(vk::Result::ERROR_OUT_OF_DATE_KHR) => { + self.recreate_swapchain( + self.size.width as u32, + self.size.height as u32, + ); + return None; + } + Err(e) => panic!("acquire_next_image failed: {e:?}"), + } + }; + if suboptimal { + self.needs_recreate = true; + } + + // Only reset *after* we've committed to submitting — resetting + // before acquire_next_image would leave us deadlocked if the + // acquire returned OUT_OF_DATE and we bailed out. + unsafe { + self.device + .reset_fences(&[sync.in_flight]) + .expect("reset_fences"); + self.device + .reset_command_pool(sync.cmd_pool, vk::CommandPoolResetFlags::empty()) + .expect("reset_command_pool"); + self.device + .begin_command_buffer( + sync.cmd_buffer, + &vk::CommandBufferBeginInfo::default() + .flags(vk::CommandBufferUsageFlags::ONE_TIME_SUBMIT), + ) + .expect("begin_command_buffer"); + } + + Some(VulkanFrame { + image_index, + image: self.swapchain_images[image_index as usize], + image_view: self.swapchain_views[image_index as usize], + cmd_buffer: sync.cmd_buffer, + extent: self.swapchain_extent, + format: self.swapchain_format, + slot, + }) + } + + /// End the command buffer, submit, present, advance frame index. + pub fn present_frame(&mut self, frame: VulkanFrame) { + let sync = &self.frames[frame.slot]; + unsafe { + self.device + .end_command_buffer(sync.cmd_buffer) + .expect("end_command_buffer"); + + let wait_semaphores = [sync.image_available]; + let wait_stages = [vk::PipelineStageFlags::COLOR_ATTACHMENT_OUTPUT]; + let signal_semaphores = [sync.render_finished]; + let cmd_buffers = [sync.cmd_buffer]; + let submit = vk::SubmitInfo::default() + .wait_semaphores(&wait_semaphores) + .wait_dst_stage_mask(&wait_stages) + .command_buffers(&cmd_buffers) + .signal_semaphores(&signal_semaphores); + self.device + .queue_submit(self.queue, &[submit], sync.in_flight) + .expect("queue_submit"); + + let swapchains = [self.swapchain]; + let image_indices = [frame.image_index]; + let present_info = vk::PresentInfoKHR::default() + .wait_semaphores(&signal_semaphores) + .swapchains(&swapchains) + .image_indices(&image_indices); + match self + .swapchain_loader + .queue_present(self.queue, &present_info) + { + Ok(suboptimal) => { + if suboptimal { + self.needs_recreate = true; + } + } + Err(vk::Result::ERROR_OUT_OF_DATE_KHR) => { + self.needs_recreate = true; + } + Err(e) => panic!("queue_present failed: {e:?}"), + } + } + + self.frame_index = (self.frame_index + 1) % FRAMES_IN_FLIGHT; + } + + /// Expose the underlying device so the renderer can record commands. + #[inline] + pub fn device(&self) -> &Device { + &self.device + } + + /// Color attachment format the swapchain was created with. Real + /// pipelines need this at construction time so `VkPipelineRenderingCreateInfo` + /// can declare a matching color attachment format. Stable across + /// resize (only `extent` changes there). + #[inline] + pub fn swapchain_format(&self) -> vk::Format { + self.swapchain_format + } + + /// The instance + physical device that own this context's logical + /// device. Renderers cache these so they can allocate buffers / + /// images via the free `allocate_host_visible_buffer_raw` / + /// `allocate_sampled_image_raw` helpers without needing a live + /// `&VulkanContext` borrow at every allocation site (chiefly, + /// `resize` which only has `&mut self`). + #[inline] + pub fn instance(&self) -> &Instance { + &self.instance + } + + #[inline] + pub fn physical_device(&self) -> vk::PhysicalDevice { + self.physical_device + } + + /// Pipeline cache shared by every renderer's + /// `create_graphics_pipelines` call. Pass this instead of + /// `vk::PipelineCache::null()` so cached binaries land on disk + /// at shutdown and short-circuit subsequent compiles. + #[inline] + pub fn pipeline_cache(&self) -> vk::PipelineCache { + self.pipeline_cache + } + + /// Run `record` against a transient command buffer, submit it, + /// wait for completion. Used for one-shot transfer work that + /// can't piggy-back on the per-frame command buffer (atlas / + /// image / texture uploads triggered from outside the render + /// loop, where there's no live `cmd` to append to). + /// + /// Allocates a fresh `vk::CommandPool` + `vk::Fence` per call + /// and tears them down at the end. Cheap (microseconds) compared + /// to the actual GPU transfer; not a hot path. + pub fn submit_oneshot(&self, record: F) { + unsafe { + let pool_info = vk::CommandPoolCreateInfo::default() + .queue_family_index(self.queue_family_index) + .flags(vk::CommandPoolCreateFlags::TRANSIENT); + let pool = self + .device + .create_command_pool(&pool_info, None) + .expect("create_command_pool(oneshot)"); + + let alloc = vk::CommandBufferAllocateInfo::default() + .command_pool(pool) + .level(vk::CommandBufferLevel::PRIMARY) + .command_buffer_count(1); + let cmd = self + .device + .allocate_command_buffers(&alloc) + .expect("allocate_command_buffers(oneshot)")[0]; + + let begin = vk::CommandBufferBeginInfo::default() + .flags(vk::CommandBufferUsageFlags::ONE_TIME_SUBMIT); + self.device + .begin_command_buffer(cmd, &begin) + .expect("begin_command_buffer(oneshot)"); + + record(cmd); + + self.device + .end_command_buffer(cmd) + .expect("end_command_buffer(oneshot)"); + + let fence = self + .device + .create_fence(&vk::FenceCreateInfo::default(), None) + .expect("create_fence(oneshot)"); + let cmds = [cmd]; + let submit = vk::SubmitInfo::default().command_buffers(&cmds); + self.device + .queue_submit(self.queue, &[submit], fence) + .expect("queue_submit(oneshot)"); + self.device + .wait_for_fences(&[fence], true, u64::MAX) + .expect("wait_for_fences(oneshot)"); + + self.device.destroy_fence(fence, None); + self.device.destroy_command_pool(pool, None); + } + } + + /// Index of the slot the *next* `acquire_frame` will use. Renderers + /// (grid, text, image overlay) ring their per-frame GPU resources by + /// this index so that a write into slot N can't race the GPU still + /// reading from slot N. Stable for the lifetime of `VulkanContext`. + #[inline] + pub fn current_frame_slot(&self) -> usize { + self.frame_index + } + + /// Allocate a host-visible, host-coherent, persistently-mapped buffer + /// suitable for per-frame uploads (uniform buffers, vertex/instance + /// buffers, storage buffers that the CPU writes into and the GPU + /// reads from this frame). On UMA/integrated GPUs the underlying + /// memory will also be `DEVICE_LOCAL` (BAR memory) — the driver + /// picks the best matching type via `memory_type_bits` filtering. + /// + /// We do not run a suballocator: each call burns one device memory + /// allocation. Vulkan guarantees ≥4096 active allocations per + /// device, and sugarloaf's working set is well under that ceiling + /// (a couple of atlases + per-frame ring buffers per terminal). + /// Switch to a slab allocator only if profiling ever shows + /// allocation churn — current call sites construct once, reuse + /// thereafter, and only reallocate on grow. + pub fn allocate_host_visible_buffer( + &self, + size: u64, + usage: vk::BufferUsageFlags, + ) -> VulkanBuffer { + // `vkCreateBuffer` rejects zero-sized buffers; bump up to a + // single byte so callers don't have to special-case empty rings. + let size = size.max(1); + + let buffer_info = vk::BufferCreateInfo::default() + .size(size) + .usage(usage) + .sharing_mode(vk::SharingMode::EXCLUSIVE); + let buffer = unsafe { + self.device + .create_buffer(&buffer_info, None) + .expect("create_buffer") + }; + + let req = unsafe { self.device.get_buffer_memory_requirements(buffer) }; + let mem_type = find_memory_type( + &self.instance, + self.physical_device, + req.memory_type_bits, + vk::MemoryPropertyFlags::HOST_VISIBLE + | vk::MemoryPropertyFlags::HOST_COHERENT, + ) + .expect("no HOST_VISIBLE | HOST_COHERENT memory type — driver bug?"); + + let alloc_info = vk::MemoryAllocateInfo::default() + .allocation_size(req.size) + .memory_type_index(mem_type); + let memory = unsafe { + self.device + .allocate_memory(&alloc_info, None) + .expect("allocate_memory") + }; + unsafe { + self.device + .bind_buffer_memory(buffer, memory, 0) + .expect("bind_buffer_memory"); + } + + // HOST_COHERENT means we never have to flush; mapping stays + // valid until `vkUnmapMemory`, which we only do at Drop. + let mapped = unsafe { + self.device + .map_memory(memory, 0, vk::WHOLE_SIZE, vk::MemoryMapFlags::empty()) + .expect("map_memory") as *mut u8 + }; + + VulkanBuffer { + device: self.device.clone(), + buffer, + memory, + mapped, + size, + } + } +} + +/// Host-visible, persistently-mapped buffer. Written to via [`as_mut_ptr`] +/// (raw pointer; caller owns the layout / bounds checks). The buffer +/// destroys itself + frees its backing memory on drop. +pub struct VulkanBuffer { + device: Device, + buffer: vk::Buffer, + memory: vk::DeviceMemory, + mapped: *mut u8, + size: u64, +} + +// `vk::Buffer`, `vk::DeviceMemory`, and the mapped pointer are all +// values the driver hands out per-allocation; ash::Device is itself +// Send+Sync. Buffers are never shared across threads in sugarloaf, but +// `Send` lets them sit inside `Sugarloaf` (which is not `!Send`). +unsafe impl Send for VulkanBuffer {} +unsafe impl Sync for VulkanBuffer {} + +impl VulkanBuffer { + #[inline] + pub fn handle(&self) -> vk::Buffer { + self.buffer + } + + #[inline] + pub fn size(&self) -> u64 { + self.size + } + + /// Raw pointer to the start of the mapping. Valid for the lifetime + /// of this `VulkanBuffer`. Writes through this pointer are visible + /// to the GPU at submit time — `HOST_COHERENT` removes the need for + /// `vkFlushMappedMemoryRanges`. + #[inline] + pub fn as_mut_ptr(&self) -> *mut u8 { + self.mapped + } +} + +impl Drop for VulkanBuffer { + fn drop(&mut self) { + unsafe { + // Order: unmap, free memory, destroy buffer. + // `vkFreeMemory` on a non-mapped allocation is safe; we + // unmap first only because some validation layers warn + // about freeing memory that's still mapped. + self.device.unmap_memory(self.memory); + self.device.destroy_buffer(self.buffer, None); + self.device.free_memory(self.memory, None); + } + } +} + +/// Free-function variant of [`VulkanContext::allocate_host_visible_buffer`] +/// for callers that hold cached `(device, instance, physical_device)` +/// rather than a live `&VulkanContext` borrow. The grid / text / +/// image renderers stash those handles at construction time so they +/// can allocate from inside their own `resize` paths (which only have +/// `&mut self`, not the parent context). +pub fn allocate_host_visible_buffer_raw( + device: &Device, + instance: &Instance, + physical_device: vk::PhysicalDevice, + size: u64, + usage: vk::BufferUsageFlags, +) -> VulkanBuffer { + let size = size.max(1); + let buffer_info = vk::BufferCreateInfo::default() + .size(size) + .usage(usage) + .sharing_mode(vk::SharingMode::EXCLUSIVE); + let buffer = unsafe { + device + .create_buffer(&buffer_info, None) + .expect("create_buffer") + }; + let req = unsafe { device.get_buffer_memory_requirements(buffer) }; + let mem_type = find_memory_type( + instance, + physical_device, + req.memory_type_bits, + vk::MemoryPropertyFlags::HOST_VISIBLE | vk::MemoryPropertyFlags::HOST_COHERENT, + ) + .expect("no HOST_VISIBLE | HOST_COHERENT memory type"); + let alloc_info = vk::MemoryAllocateInfo::default() + .allocation_size(req.size) + .memory_type_index(mem_type); + let memory = unsafe { + device + .allocate_memory(&alloc_info, None) + .expect("allocate_memory") + }; + unsafe { + device + .bind_buffer_memory(buffer, memory, 0) + .expect("bind_buffer_memory"); + } + let mapped = unsafe { + device + .map_memory(memory, 0, vk::WHOLE_SIZE, vk::MemoryMapFlags::empty()) + .expect("map_memory") as *mut u8 + }; + VulkanBuffer { + device: device.clone(), + buffer, + memory, + mapped, + size, + } +} + +/// Walks the device's memory types looking for one that matches both +/// `type_filter` (the bitmask returned by `vkGetBufferMemoryRequirements`) +/// and the requested `flags`. Returns `None` if no matching type +/// exists — that's a Vulkan-spec violation the driver should never +/// produce for the standard `HOST_VISIBLE | HOST_COHERENT` and +/// `DEVICE_LOCAL` combinations, but callers should still treat it as +/// fatal rather than ignore it. +fn find_memory_type( + instance: &Instance, + physical_device: vk::PhysicalDevice, + type_filter: u32, + flags: vk::MemoryPropertyFlags, +) -> Option { + let props = + unsafe { instance.get_physical_device_memory_properties(physical_device) }; + for i in 0..props.memory_type_count { + let supported = (type_filter & (1 << i)) != 0; + let matches_flags = props.memory_types[i as usize] + .property_flags + .contains(flags); + if supported && matches_flags { + return Some(i); + } + } + None +} + +/// One-off boot-time log of the memory types we'd pick for our two +/// hot allocation patterns. On UMA / integrated GPUs (Intel iGPU, +/// AMD APU, common Debian-laptop hardware) we expect the +/// host-visible heap to also report `DEVICE_LOCAL` — that's BAR +/// memory and our persistently-mapped per-frame buffers land in fast +/// GPU-accessible memory with no staging copy. On discrete GPUs the +/// host-visible heap is plain system RAM, slower for the GPU to +/// read; we'd want to switch to staging-buffer uploads for hot +/// per-frame data if profiling shows it matters. +fn log_memory_heap_choice(instance: &Instance, physical_device: vk::PhysicalDevice) { + // Pretend `type_filter = !0` to ignore per-resource alignment + // filtering — we just want the canonical pick for each pattern. + let host_visible = find_memory_type( + instance, + physical_device, + !0, + vk::MemoryPropertyFlags::HOST_VISIBLE | vk::MemoryPropertyFlags::HOST_COHERENT, + ); + let device_local = find_memory_type( + instance, + physical_device, + !0, + vk::MemoryPropertyFlags::DEVICE_LOCAL, + ); + let props = + unsafe { instance.get_physical_device_memory_properties(physical_device) }; + if let Some(idx) = host_visible { + let flags = props.memory_types[idx as usize].property_flags; + let bar = flags.contains(vk::MemoryPropertyFlags::DEVICE_LOCAL); + tracing::info!( + "Vulkan host-visible memory: type {} flags={:?} ({})", + idx, + flags, + if bar { + "BAR / unified — fast GPU reads" + } else { + "system RAM — slower GPU reads, consider staging for hot data" + } + ); + } + if let Some(idx) = device_local { + tracing::info!( + "Vulkan device-local memory: type {} flags={:?}", + idx, + props.memory_types[idx as usize].property_flags + ); + } +} + +// ----------------------------------------------------------------------- +// Image helper (device-local 2D image + view + memory) +// ----------------------------------------------------------------------- + +impl VulkanContext { + /// Allocate a device-local 2D image suitable for sampling from a + /// shader (atlas, kitty graphic, background image). Created in + /// `UNDEFINED` layout — the caller's first transfer command must + /// include a barrier transitioning to `TRANSFER_DST_OPTIMAL` + /// before any `vkCmdCopyBufferToImage`. + pub fn allocate_sampled_image( + &self, + width: u32, + height: u32, + format: vk::Format, + usage: vk::ImageUsageFlags, + ) -> VulkanImage { + let image_info = vk::ImageCreateInfo::default() + .image_type(vk::ImageType::TYPE_2D) + .format(format) + .extent(vk::Extent3D { + width, + height, + depth: 1, + }) + .mip_levels(1) + .array_layers(1) + .samples(vk::SampleCountFlags::TYPE_1) + .tiling(vk::ImageTiling::OPTIMAL) + .usage(usage) + .sharing_mode(vk::SharingMode::EXCLUSIVE) + .initial_layout(vk::ImageLayout::UNDEFINED); + let image = unsafe { + self.device + .create_image(&image_info, None) + .expect("create_image") + }; + + let req = unsafe { self.device.get_image_memory_requirements(image) }; + let mem_type = find_memory_type( + &self.instance, + self.physical_device, + req.memory_type_bits, + vk::MemoryPropertyFlags::DEVICE_LOCAL, + ) + .expect("no DEVICE_LOCAL memory type — driver bug?"); + + let alloc_info = vk::MemoryAllocateInfo::default() + .allocation_size(req.size) + .memory_type_index(mem_type); + let memory = unsafe { + self.device + .allocate_memory(&alloc_info, None) + .expect("allocate_memory(image)") + }; + unsafe { + self.device + .bind_image_memory(image, memory, 0) + .expect("bind_image_memory"); + } + + let view_info = vk::ImageViewCreateInfo::default() + .image(image) + .view_type(vk::ImageViewType::TYPE_2D) + .format(format) + .components(vk::ComponentMapping::default()) + .subresource_range( + vk::ImageSubresourceRange::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .base_mip_level(0) + .level_count(1) + .base_array_layer(0) + .layer_count(1), + ); + let view = unsafe { + self.device + .create_image_view(&view_info, None) + .expect("create_image_view") + }; + + VulkanImage { + device: self.device.clone(), + image, + view, + memory, + width, + height, + format, + } + } +} + +/// Device-local 2D image with view + backing memory. Drops the view, +/// image, and memory on `Drop`. The image starts in `UNDEFINED` layout +/// — the first command that uses it must barrier-transition to a +/// usable layout (`TRANSFER_DST_OPTIMAL` for the initial upload). +pub struct VulkanImage { + device: Device, + image: vk::Image, + view: vk::ImageView, + memory: vk::DeviceMemory, + pub width: u32, + pub height: u32, + pub format: vk::Format, +} + +unsafe impl Send for VulkanImage {} +unsafe impl Sync for VulkanImage {} + +impl VulkanImage { + #[inline] + pub fn handle(&self) -> vk::Image { + self.image + } + + #[inline] + pub fn view(&self) -> vk::ImageView { + self.view + } +} + +impl Drop for VulkanImage { + fn drop(&mut self) { + unsafe { + self.device.destroy_image_view(self.view, None); + self.device.destroy_image(self.image, None); + self.device.free_memory(self.memory, None); + } + } +} + +impl Drop for VulkanContext { + fn drop(&mut self) { + unsafe { + let _ = self.device.device_wait_idle(); + + // Best-effort: serialize the pipeline cache to disk + // before destroying it. Failure (no XDG_CACHE_HOME, no + // write perms, etc) is logged but not fatal. + save_pipeline_cache(&self.device, self.pipeline_cache); + self.device + .destroy_pipeline_cache(self.pipeline_cache, None); + + for frame in &self.frames { + self.device.destroy_semaphore(frame.image_available, None); + self.device.destroy_semaphore(frame.render_finished, None); + self.device.destroy_fence(frame.in_flight, None); + self.device.destroy_command_pool(frame.cmd_pool, None); + } + + for &view in &self.swapchain_views { + self.device.destroy_image_view(view, None); + } + self.swapchain_loader + .destroy_swapchain(self.swapchain, None); + self.device.destroy_device(None); + self.surface_loader.destroy_surface(self.surface, None); + self.instance.destroy_instance(None); + } + } +} + +// ======================================================================= +// Pipeline cache (load on `new`, save on `Drop`) +// ======================================================================= + +/// Path to the on-disk pipeline cache. Returns `None` if neither +/// `XDG_CACHE_HOME` nor `HOME` is set. +fn pipeline_cache_path() -> Option { + let dir = if let Some(xdg) = std::env::var_os("XDG_CACHE_HOME") { + std::path::PathBuf::from(xdg) + } else if let Some(home) = std::env::var_os("HOME") { + let mut p = std::path::PathBuf::from(home); + p.push(".cache"); + p + } else { + return None; + }; + Some(dir.join("rio").join("sugarloaf-vulkan.cache")) +} + +fn create_pipeline_cache(device: &Device) -> vk::PipelineCache { + let initial_data: Vec = pipeline_cache_path() + .and_then(|p| std::fs::read(&p).ok()) + .unwrap_or_default(); + if !initial_data.is_empty() { + tracing::info!("loaded Vulkan pipeline cache: {} bytes", initial_data.len()); + } + let info = vk::PipelineCacheCreateInfo::default().initial_data(&initial_data); + unsafe { + device + .create_pipeline_cache(&info, None) + .expect("create_pipeline_cache") + } +} + +fn save_pipeline_cache(device: &Device, cache: vk::PipelineCache) { + let Some(path) = pipeline_cache_path() else { + return; + }; + let data = match unsafe { device.get_pipeline_cache_data(cache) } { + Ok(d) => d, + Err(e) => { + tracing::warn!("get_pipeline_cache_data failed: {:?}", e); + return; + } + }; + if data.is_empty() { + return; + } + if let Some(parent) = path.parent() { + if let Err(e) = std::fs::create_dir_all(parent) { + tracing::warn!("pipeline cache mkdir {:?} failed: {}", parent, e); + return; + } + } + if let Err(e) = std::fs::write(&path, &data) { + tracing::warn!("pipeline cache write {:?} failed: {}", path, e); + } else { + tracing::info!( + "saved Vulkan pipeline cache: {} bytes → {:?}", + data.len(), + path + ); + } +} + +// ------------------------------------------------------------------------- +// Internal helpers (free functions so `new()` stays readable). +// ------------------------------------------------------------------------- + +fn create_instance( + entry: &Entry, + window: &SugarloafWindow, + enable_validation: bool, +) -> Instance { + let app_name = c"sugarloaf"; + let app_info = vk::ApplicationInfo::default() + .application_name(app_name) + .application_version(0) + .engine_name(app_name) + .engine_version(0) + .api_version(vk::API_VERSION_1_3); + + // KHR_surface + the right platform surface extension for the window + // handle we were given. Adding extensions the driver doesn't + // advertise makes `create_instance` fail, so we match the window + // type exactly instead of asking for all three. + let mut extensions: Vec<*const c_char> = vec![khr::surface::NAME.as_ptr()]; + match window.display_handle().unwrap().as_raw() { + RawDisplayHandle::Xlib(_) => extensions.push(khr::xlib_surface::NAME.as_ptr()), + RawDisplayHandle::Xcb(_) => extensions.push(khr::xcb_surface::NAME.as_ptr()), + RawDisplayHandle::Wayland(_) => { + extensions.push(khr::wayland_surface::NAME.as_ptr()) + } + other => panic!("Vulkan backend: unsupported display handle {:?}", other), + } + + // Validation: append `VK_EXT_debug_utils` so we can install a + // messenger callback after instance creation. The layer + // (`VK_LAYER_KHRONOS_validation`) is enabled separately via + // `enabled_layer_names` below. + let validation_layer_name = c"VK_LAYER_KHRONOS_validation"; + let layer_ptrs: Vec<*const c_char> = if enable_validation { + if validation_layer_available(entry, validation_layer_name) { + extensions.push(ash::ext::debug_utils::NAME.as_ptr()); + vec![validation_layer_name.as_ptr()] + } else { + tracing::warn!( + "RIO_VULKAN_VALIDATION set but VK_LAYER_KHRONOS_validation \ + not available — install `vulkan-validationlayers` (Debian) \ + / `vulkan-validation-layers` (Arch) to enable it" + ); + Vec::new() + } + } else { + Vec::new() + }; + + let create_info = vk::InstanceCreateInfo::default() + .application_info(&app_info) + .enabled_extension_names(&extensions) + .enabled_layer_names(&layer_ptrs); + + unsafe { entry.create_instance(&create_info, None) } + .expect("vkCreateInstance failed — is a Vulkan 1.3 driver installed?") +} + +/// True if the user opted into validation via `RIO_VULKAN_VALIDATION=1`. +/// We always read the env var (debug + release) so users can flip it +/// on for one run without recompiling. +fn validation_requested() -> bool { + std::env::var_os("RIO_VULKAN_VALIDATION") + .map(|v| v != "0" && !v.is_empty()) + .unwrap_or(false) +} + +fn validation_layer_available(entry: &Entry, target: &CStr) -> bool { + match unsafe { entry.enumerate_instance_layer_properties() } { + Ok(layers) => layers.iter().any(|l| { + let name = unsafe { CStr::from_ptr(l.layer_name.as_ptr()) }; + name == target + }), + Err(_) => false, + } +} + +fn create_debug_messenger(entry: &Entry, instance: &Instance) -> Option { + let loader = ash::ext::debug_utils::Instance::new(entry, instance); + + let info = vk::DebugUtilsMessengerCreateInfoEXT::default() + .message_severity( + vk::DebugUtilsMessageSeverityFlagsEXT::ERROR + | vk::DebugUtilsMessageSeverityFlagsEXT::WARNING + | vk::DebugUtilsMessageSeverityFlagsEXT::INFO, + ) + .message_type( + vk::DebugUtilsMessageTypeFlagsEXT::GENERAL + | vk::DebugUtilsMessageTypeFlagsEXT::VALIDATION + | vk::DebugUtilsMessageTypeFlagsEXT::PERFORMANCE, + ) + .pfn_user_callback(Some(debug_callback)); + + let handle = unsafe { loader.create_debug_utils_messenger(&info, None) } + .expect("create_debug_utils_messenger"); + tracing::info!("Vulkan validation layers active"); + Some(DebugMessenger { loader, handle }) +} + +unsafe extern "system" fn debug_callback( + severity: vk::DebugUtilsMessageSeverityFlagsEXT, + msg_type: vk::DebugUtilsMessageTypeFlagsEXT, + callback_data: *const vk::DebugUtilsMessengerCallbackDataEXT<'_>, + _user_data: *mut std::ffi::c_void, +) -> vk::Bool32 { + let data = unsafe { &*callback_data }; + let message = if data.p_message.is_null() { + std::borrow::Cow::Borrowed("") + } else { + unsafe { CStr::from_ptr(data.p_message) }.to_string_lossy() + }; + let kind = if msg_type.contains(vk::DebugUtilsMessageTypeFlagsEXT::VALIDATION) { + "validation" + } else if msg_type.contains(vk::DebugUtilsMessageTypeFlagsEXT::PERFORMANCE) { + "perf" + } else { + "general" + }; + if severity.contains(vk::DebugUtilsMessageSeverityFlagsEXT::ERROR) { + tracing::error!("vk[{}] {}", kind, message); + } else if severity.contains(vk::DebugUtilsMessageSeverityFlagsEXT::WARNING) { + tracing::warn!("vk[{}] {}", kind, message); + } else if severity.contains(vk::DebugUtilsMessageSeverityFlagsEXT::INFO) { + tracing::info!("vk[{}] {}", kind, message); + } else { + tracing::debug!("vk[{}] {}", kind, message); + } + vk::FALSE +} + +fn create_surface( + entry: &Entry, + instance: &Instance, + window: &SugarloafWindow, +) -> vk::SurfaceKHR { + let display = window.display_handle().unwrap().as_raw(); + let window_handle = window.window_handle().unwrap().as_raw(); + + unsafe { + match (display, window_handle) { + (RawDisplayHandle::Xlib(d), RawWindowHandle::Xlib(w)) => { + let loader = khr::xlib_surface::Instance::new(entry, instance); + let info = vk::XlibSurfaceCreateInfoKHR::default() + .dpy( + d.display + .expect("Xlib display pointer missing") + .as_ptr() + .cast(), + ) + .window(w.window); + loader + .create_xlib_surface(&info, None) + .expect("create_xlib_surface") + } + (RawDisplayHandle::Xcb(d), RawWindowHandle::Xcb(w)) => { + let loader = khr::xcb_surface::Instance::new(entry, instance); + let info = vk::XcbSurfaceCreateInfoKHR::default() + .connection( + d.connection + .expect("Xcb connection pointer missing") + .as_ptr() + .cast(), + ) + .window(w.window.get()); + loader + .create_xcb_surface(&info, None) + .expect("create_xcb_surface") + } + (RawDisplayHandle::Wayland(d), RawWindowHandle::Wayland(w)) => { + let loader = khr::wayland_surface::Instance::new(entry, instance); + let info = vk::WaylandSurfaceCreateInfoKHR::default() + .display(d.display.as_ptr().cast()) + .surface(w.surface.as_ptr().cast()); + loader + .create_wayland_surface(&info, None) + .expect("create_wayland_surface") + } + (d, w) => panic!( + "Vulkan backend: mismatched or unsupported handles: display={d:?} window={w:?}" + ), + } + } +} + +/// Pick a physical device + queue family. Prefer discrete GPU, require +/// a queue family that supports both graphics and present on our surface. +fn pick_physical_device( + instance: &Instance, + surface_loader: &khr::surface::Instance, + surface: vk::SurfaceKHR, +) -> (vk::PhysicalDevice, u32) { + let devices = unsafe { instance.enumerate_physical_devices() } + .expect("enumerate_physical_devices"); + + let mut best: Option<(vk::PhysicalDevice, u32, i32)> = None; + for device in devices { + let props = unsafe { instance.get_physical_device_properties(device) }; + let qf_props = + unsafe { instance.get_physical_device_queue_family_properties(device) }; + + for (index, qf) in qf_props.iter().enumerate() { + let index = index as u32; + if !qf.queue_flags.contains(vk::QueueFlags::GRAPHICS) { + continue; + } + let present_ok = unsafe { + surface_loader + .get_physical_device_surface_support(device, index, surface) + .unwrap_or(false) + }; + if !present_ok { + continue; + } + let score = match props.device_type { + vk::PhysicalDeviceType::DISCRETE_GPU => 1000, + vk::PhysicalDeviceType::INTEGRATED_GPU => 500, + vk::PhysicalDeviceType::VIRTUAL_GPU => 100, + vk::PhysicalDeviceType::CPU => 10, + _ => 1, + }; + if best.map(|(_, _, s)| score > s).unwrap_or(true) { + best = Some((device, index, score)); + } + } + } + + let (device, queue_family, _) = + best.expect("no Vulkan device with graphics + present support on this surface"); + (device, queue_family) +} + +fn physical_device_name(instance: &Instance, device: vk::PhysicalDevice) -> String { + let props = unsafe { instance.get_physical_device_properties(device) }; + // `device_name` is a C string embedded in a fixed-size array. + let raw = props.device_name.as_ptr(); + unsafe { CStr::from_ptr(raw) } + .to_string_lossy() + .into_owned() +} + +fn create_device( + instance: &Instance, + physical_device: vk::PhysicalDevice, + queue_family_index: u32, +) -> Device { + let queue_priorities = [1.0f32]; + let queue_info = vk::DeviceQueueCreateInfo::default() + .queue_family_index(queue_family_index) + .queue_priorities(&queue_priorities); + + let device_extensions = [khr::swapchain::NAME.as_ptr()]; + + // Enable dynamic rendering up front — it's Vulkan 1.3 core. We're + // not using it yet in the clear-only path, but enabling it here + // avoids having to recreate the device when pipelines land. + let mut vk13_features = + vk::PhysicalDeviceVulkan13Features::default().dynamic_rendering(true); + + let queue_infos = [queue_info]; + let create_info = vk::DeviceCreateInfo::default() + .queue_create_infos(&queue_infos) + .enabled_extension_names(&device_extensions) + .push_next(&mut vk13_features); + + unsafe { instance.create_device(physical_device, &create_info, None) } + .expect("vkCreateDevice") +} + +/// Build a swapchain and its image views. `old` is passed as +/// `old_swapchain` so the driver can recycle images during resize. +#[allow(clippy::too_many_arguments)] +fn create_swapchain( + device: &Device, + surface_loader: &khr::surface::Instance, + swapchain_loader: &khr::swapchain::Device, + physical_device: vk::PhysicalDevice, + surface: vk::SurfaceKHR, + requested_width: u32, + requested_height: u32, + old: vk::SwapchainKHR, +) -> ( + vk::SwapchainKHR, + vk::Format, + vk::ColorSpaceKHR, + vk::Extent2D, + Vec, + Vec, +) { + let caps = unsafe { + surface_loader + .get_physical_device_surface_capabilities(physical_device, surface) + .expect("get_physical_device_surface_capabilities") + }; + let formats = unsafe { + surface_loader + .get_physical_device_surface_formats(physical_device, surface) + .expect("get_physical_device_surface_formats") + }; + let present_modes = unsafe { + surface_loader + .get_physical_device_surface_present_modes(physical_device, surface) + .expect("get_physical_device_surface_present_modes") + }; + + // Prefer BGRA8_UNORM (linear) so blending stays in gamma space — the + // same choice Metal makes (`MTLPixelFormat::BGRA8Unorm` + DisplayP3 + // tag). Fragment shaders will emit sRGB-encoded output. If the + // driver doesn't offer BGRA8_UNORM, fall back to whatever it gives + // us — formats[0] is guaranteed present per the spec. + let chosen_format = formats + .iter() + .find(|f| { + f.format == vk::Format::B8G8R8A8_UNORM + && f.color_space == vk::ColorSpaceKHR::SRGB_NONLINEAR + }) + .copied() + .unwrap_or(formats[0]); + + // Present mode: Mailbox for low-latency "triple buffered", FIFO as a + // guaranteed fallback. We don't expose a config knob yet — same + // story as Metal's hard-coded `maximumDrawableCount = 3`. + let present_mode = if present_modes.contains(&vk::PresentModeKHR::MAILBOX) { + vk::PresentModeKHR::MAILBOX + } else { + vk::PresentModeKHR::FIFO + }; + + let extent = if caps.current_extent.width != u32::MAX { + caps.current_extent + } else { + vk::Extent2D { + width: requested_width + .clamp(caps.min_image_extent.width, caps.max_image_extent.width), + height: requested_height + .clamp(caps.min_image_extent.height, caps.max_image_extent.height), + } + }; + + // Aim for 3 images where the driver allows it (triple buffering), + // clamped to the advertised range. `max_image_count == 0` means "no + // upper limit". + let mut image_count = caps.min_image_count.max(3); + if caps.max_image_count != 0 && image_count > caps.max_image_count { + image_count = caps.max_image_count; + } + + let create_info = vk::SwapchainCreateInfoKHR::default() + .surface(surface) + .min_image_count(image_count) + .image_format(chosen_format.format) + .image_color_space(chosen_format.color_space) + .image_extent(extent) + .image_array_layers(1) + .image_usage( + vk::ImageUsageFlags::COLOR_ATTACHMENT | vk::ImageUsageFlags::TRANSFER_DST, + ) + .image_sharing_mode(vk::SharingMode::EXCLUSIVE) + .pre_transform(caps.current_transform) + .composite_alpha(vk::CompositeAlphaFlagsKHR::OPAQUE) + .present_mode(present_mode) + .clipped(true) + .old_swapchain(old); + + let swapchain = unsafe { swapchain_loader.create_swapchain(&create_info, None) } + .expect("create_swapchain"); + + let images = unsafe { swapchain_loader.get_swapchain_images(swapchain) } + .expect("get_swapchain_images"); + + let views = images + .iter() + .map(|&image| { + let info = vk::ImageViewCreateInfo::default() + .image(image) + .view_type(vk::ImageViewType::TYPE_2D) + .format(chosen_format.format) + .components(vk::ComponentMapping::default()) + .subresource_range( + vk::ImageSubresourceRange::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .base_mip_level(0) + .level_count(1) + .base_array_layer(0) + .layer_count(1), + ); + unsafe { device.create_image_view(&info, None) }.expect("create_image_view") + }) + .collect(); + + ( + swapchain, + chosen_format.format, + chosen_format.color_space, + extent, + images, + views, + ) +} + +fn create_frames( + device: &Device, + queue_family_index: u32, +) -> [FrameSync; FRAMES_IN_FLIGHT] { + std::array::from_fn(|_| { + let semaphore_info = vk::SemaphoreCreateInfo::default(); + let fence_info = + vk::FenceCreateInfo::default().flags(vk::FenceCreateFlags::SIGNALED); + let pool_info = vk::CommandPoolCreateInfo::default() + .queue_family_index(queue_family_index) + .flags(vk::CommandPoolCreateFlags::TRANSIENT); + + unsafe { + let image_available = device + .create_semaphore(&semaphore_info, None) + .expect("create_semaphore"); + let render_finished = device + .create_semaphore(&semaphore_info, None) + .expect("create_semaphore"); + let in_flight = device + .create_fence(&fence_info, None) + .expect("create_fence"); + let cmd_pool = device + .create_command_pool(&pool_info, None) + .expect("create_command_pool"); + let alloc_info = vk::CommandBufferAllocateInfo::default() + .command_pool(cmd_pool) + .level(vk::CommandBufferLevel::PRIMARY) + .command_buffer_count(1); + let cmd_buffer = device + .allocate_command_buffers(&alloc_info) + .expect("allocate_command_buffers")[0]; + + FrameSync { + image_available, + render_finished, + in_flight, + cmd_pool, + cmd_buffer, + } + } + }) +} diff --git a/sugarloaf/src/context/webgpu.rs b/sugarloaf/src/context/webgpu.rs index 794b3a6e..aa0ed944 100644 --- a/sugarloaf/src/context/webgpu.rs +++ b/sugarloaf/src/context/webgpu.rs @@ -59,7 +59,12 @@ impl<'a> WgpuContext<'a> { instance.create_surface(sugarloaf_window).unwrap(); let adapter = futures::executor::block_on(instance.request_adapter( &wgpu::RequestAdapterOptions { - power_preference: renderer_config.power_preference, + // Hard-coded — sugarloaf used to expose a + // `power_preference` knob, but in practice every Rio + // user picks `HighPerformance` (the alternative gives + // visibly worse text on hybrid laptops). Removed from + // the public API. + power_preference: wgpu::PowerPreference::HighPerformance, compatible_surface: Some(&surface), force_fallback_adapter: false, }, diff --git a/sugarloaf/src/grid/cell.rs b/sugarloaf/src/grid/cell.rs index f883175c..a6854b0e 100644 --- a/sugarloaf/src/grid/cell.rs +++ b/sugarloaf/src/grid/cell.rs @@ -174,5 +174,5 @@ impl GridUniforms { const _: () = { // Keep the uniform block a multiple of 16 bytes (WGSL / std140). - assert!(std::mem::size_of::() % 16 == 0); + assert!(std::mem::size_of::().is_multiple_of(16)); }; diff --git a/sugarloaf/src/grid/mod.rs b/sugarloaf/src/grid/mod.rs index 520bac0b..073aba34 100644 --- a/sugarloaf/src/grid/mod.rs +++ b/sugarloaf/src/grid/mod.rs @@ -23,6 +23,9 @@ pub mod atlas; pub mod cell; #[cfg(target_os = "macos")] pub mod metal; +#[cfg(target_os = "linux")] +pub mod vulkan; +#[cfg(feature = "wgpu")] pub mod webgpu; use crate::context::{Context, ContextType}; @@ -37,10 +40,22 @@ pub use cell::{CellBg, CellText, GridUniforms}; /// Backend selection matches sugarloaf's existing `ContextType` — /// there's no separate config knob for grid vs. rich-text, because /// the rich-text terminal path is being removed. +/// +/// The Vulkan variant is significantly larger than the others +/// because atlas and buffer-ring state lives inline, but only one +/// variant is ever constructed per panel; boxing the bigger one +/// would just trade stack size for one allocation per panel — not +/// worth it. +#[allow(clippy::large_enum_variant)] pub enum GridRenderer { #[cfg(target_os = "macos")] Metal(metal::MetalGridRenderer), + #[cfg(feature = "wgpu")] Wgpu(webgpu::WgpuGridRenderer), + /// Native Vulkan grid renderer. Phase 3 = bg pass; text pass + + /// atlases land in Phase 4. + #[cfg(target_os = "linux")] + Vulkan(vulkan::VulkanGridRenderer), /// CPU backend has no grid renderer — it falls back to /// rasterising via the existing cpu path. Terminal content won't /// use the grid path on CPU builds. @@ -57,9 +72,16 @@ impl GridRenderer { ContextType::Metal(ctx) => { GridRenderer::Metal(metal::MetalGridRenderer::new(ctx, cols, rows)) } + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), + #[cfg(feature = "wgpu")] ContextType::Wgpu(ctx) => { GridRenderer::Wgpu(webgpu::WgpuGridRenderer::new(ctx, cols, rows)) } + #[cfg(target_os = "linux")] + ContextType::Vulkan(ctx) => { + GridRenderer::Vulkan(vulkan::VulkanGridRenderer::new(ctx, cols, rows)) + } ContextType::Cpu(_) => GridRenderer::Unsupported, } } @@ -70,7 +92,10 @@ impl GridRenderer { match self { #[cfg(target_os = "macos")] GridRenderer::Metal(r) => r.resize(cols, rows), + #[cfg(feature = "wgpu")] GridRenderer::Wgpu(r) => r.resize(cols, rows), + #[cfg(target_os = "linux")] + GridRenderer::Vulkan(r) => r.resize(cols, rows), GridRenderer::Unsupported => {} } } @@ -84,7 +109,10 @@ impl GridRenderer { match self { #[cfg(target_os = "macos")] GridRenderer::Metal(r) => r.write_row(row, bg, fg), + #[cfg(feature = "wgpu")] GridRenderer::Wgpu(r) => r.write_row(row, bg, fg), + #[cfg(target_os = "linux")] + GridRenderer::Vulkan(r) => r.write_row(row, bg, fg), GridRenderer::Unsupported => {} } } @@ -95,7 +123,10 @@ impl GridRenderer { match self { #[cfg(target_os = "macos")] GridRenderer::Metal(r) => r.clear_row(row), + #[cfg(feature = "wgpu")] GridRenderer::Wgpu(r) => r.clear_row(row), + #[cfg(target_os = "linux")] + GridRenderer::Vulkan(r) => r.clear_row(row), GridRenderer::Unsupported => {} } } @@ -119,6 +150,7 @@ impl GridRenderer { /// Wgpu counterpart of `render_metal`. Phase 1b will record a bg /// pass against the caller's `wgpu::RenderPass`. + #[cfg(feature = "wgpu")] pub fn render_wgpu<'pass>( &'pass mut self, render_pass: &mut wgpu::RenderPass<'pass>, @@ -129,6 +161,42 @@ impl GridRenderer { } } + /// 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 + /// pass. No-op for non-Vulkan renderers (Metal handles uploads + /// inside its own `replace_region`; wgpu handles them via + /// `queue.write_texture`). + #[cfg(target_os = "linux")] + pub fn prepare_vulkan( + &mut self, + ctx: &crate::context::vulkan::VulkanContext, + cmd_buffer: ash::vk::CommandBuffer, + frame_slot: usize, + ) { + if let GridRenderer::Vulkan(r) = self { + r.prepare(ctx, cmd_buffer, frame_slot); + } + } + + #[cfg(target_os = "linux")] + pub fn render_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(ctx, cmd_buffer, frame_slot, uniforms); + } + } + /// Whether this backend actually renders grid cells. Call sites /// can fall back to the rich-text path when this returns false. pub fn is_active(&self) -> bool { @@ -141,7 +209,10 @@ impl GridRenderer { match self { #[cfg(target_os = "macos")] GridRenderer::Metal(r) => r.lookup_glyph(key), + #[cfg(feature = "wgpu")] GridRenderer::Wgpu(r) => r.lookup_glyph(key), + #[cfg(target_os = "linux")] + GridRenderer::Vulkan(r) => r.lookup_glyph(key), GridRenderer::Unsupported => None, } } @@ -152,7 +223,10 @@ impl GridRenderer { match self { #[cfg(target_os = "macos")] GridRenderer::Metal(r) => r.lookup_glyph_color(key), + #[cfg(feature = "wgpu")] GridRenderer::Wgpu(r) => r.lookup_glyph_color(key), + #[cfg(target_os = "linux")] + GridRenderer::Vulkan(r) => r.lookup_glyph_color(key), GridRenderer::Unsupported => None, } } @@ -167,7 +241,10 @@ impl GridRenderer { match self { #[cfg(target_os = "macos")] GridRenderer::Metal(r) => r.insert_glyph(key, glyph), + #[cfg(feature = "wgpu")] GridRenderer::Wgpu(r) => r.insert_glyph(key, glyph), + #[cfg(target_os = "linux")] + GridRenderer::Vulkan(r) => r.insert_glyph(key, glyph), GridRenderer::Unsupported => None, } } @@ -182,7 +259,10 @@ impl GridRenderer { match self { #[cfg(target_os = "macos")] GridRenderer::Metal(r) => r.insert_glyph_color(key, glyph), + #[cfg(feature = "wgpu")] GridRenderer::Wgpu(r) => r.insert_glyph_color(key, glyph), + #[cfg(target_os = "linux")] + GridRenderer::Vulkan(r) => r.insert_glyph_color(key, glyph), GridRenderer::Unsupported => None, } } @@ -194,7 +274,10 @@ impl GridRenderer { match self { #[cfg(target_os = "macos")] GridRenderer::Metal(r) => r.needs_full_rebuild(), + #[cfg(feature = "wgpu")] GridRenderer::Wgpu(r) => r.needs_full_rebuild(), + #[cfg(target_os = "linux")] + GridRenderer::Vulkan(r) => r.needs_full_rebuild(), GridRenderer::Unsupported => false, } } @@ -205,7 +288,10 @@ impl GridRenderer { match self { #[cfg(target_os = "macos")] GridRenderer::Metal(r) => r.mark_full_rebuild_done(), + #[cfg(feature = "wgpu")] GridRenderer::Wgpu(r) => r.mark_full_rebuild_done(), + #[cfg(target_os = "linux")] + GridRenderer::Vulkan(r) => r.mark_full_rebuild_done(), GridRenderer::Unsupported => {} } } diff --git a/sugarloaf/src/grid/shaders/grid_bg.frag.glsl b/sugarloaf/src/grid/shaders/grid_bg.frag.glsl new file mode 100644 index 00000000..811d3178 --- /dev/null +++ b/sugarloaf/src/grid/shaders/grid_bg.frag.glsl @@ -0,0 +1,151 @@ +#version 450 + +// Per-cell background fragment shader. One sample per framebuffer pixel: +// look up the owning grid cell, apply `padding_extend` clamping at the +// edges, paint the cursor color where applicable, otherwise return the +// stored CellBg colour through the sRGB↔DisplayP3 chain. +// +// Ported 1:1 from `grid_bg_fragment` in +// `sugarloaf/src/grid/shaders/grid.metal`. Field order in `Uniforms` +// must match `GridUniforms` in `sugarloaf/src/grid/cell.rs` and the +// Metal `Uniforms` struct in grid.metal — std140 layout naturally +// matches the hand-packed byte layout used over there. + +// `Uniforms` mirrors `GridUniforms` (144B + pad to 160B = 10 × vec4 +// blocks under std140). See cell.rs for offsets / sizes; the fields +// here are listed in the same order. +layout(set = 0, binding = 0, std140) uniform Uniforms { + mat4 projection; // offset 0 + vec4 grid_padding; // offset 64 (top, right, bottom, left) + vec4 cursor_color; // offset 80 + vec4 cursor_bg_color; // offset 96 + vec2 cell_size; // offset 112 + uvec2 grid_size; // offset 120 + uvec2 cursor_pos; // offset 128 + uvec2 _pad_cursor; // offset 136 + float min_contrast; // offset 144 + uint flags; // offset 148 + uint padding_extend; // offset 152 + uint input_colorspace; // offset 156 +} uniforms; + +// `CellBg` is 4 bytes (uchar4 rgba) in cell.rs. We expose the storage +// buffer as `uint cells[]` and unpack each entry via +// `unpackUnorm4x8` — that maps byte[0]→x etc., matching the little- +// endian RGBA layout sugarloaf uses on the CPU side. +layout(set = 0, binding = 1, std430) readonly buffer Cells { + uint cells[]; +}; + +layout(location = 0) out vec4 out_color; + +const uint PAD_EXTEND_LEFT = 1u << 0; +const uint PAD_EXTEND_RIGHT = 1u << 1; +const uint PAD_EXTEND_UP = 1u << 2; +const uint PAD_EXTEND_DOWN = 1u << 3; + +// ----- colorspace helpers (mirror `grid.metal` 1:1) ----- + +vec3 grid_srgb_to_linear(vec3 c) { + vec3 lo = c / 12.92; + vec3 hi = pow((c + 0.055) / 1.055, vec3(2.4)); + return mix(lo, hi, greaterThan(c, vec3(0.04045))); +} + +vec3 grid_linear_to_srgb(vec3 c) { + vec3 lo = c * 12.92; + vec3 hi = pow(c, vec3(1.0 / 2.4)) * 1.055 - 0.055; + return mix(lo, hi, greaterThan(c, vec3(0.0031308))); +} + +vec3 grid_srgb_to_p3(vec3 linear_srgb) { + return vec3( + dot(linear_srgb, vec3(0.82246197, 0.17753803, 0.0)), + dot(linear_srgb, vec3(0.03319420, 0.96680580, 0.0)), + dot(linear_srgb, vec3(0.01708263, 0.07239744, 0.91051993)) + ); +} + +vec3 grid_rec2020_to_p3(vec3 linear_r2020) { + return vec3( + dot(linear_r2020, vec3( 1.34357825, -0.28217967, -0.06139858)), + dot(linear_r2020, vec3(-0.06529745, 1.08782226, -0.02252481)), + dot(linear_r2020, vec3( 0.00282179, -0.02598807, 1.02316628)) + ); +} + +vec3 grid_prepare_output_rgb(vec3 srgb, uint input_colorspace) { + vec3 lin = grid_srgb_to_linear(srgb); + if (input_colorspace == 0u) { + lin = grid_srgb_to_p3(lin); + } else if (input_colorspace == 2u) { + lin = grid_rec2020_to_p3(lin); + } + return grid_linear_to_srgb(lin); +} + +void main() { + // `gl_FragCoord.xy` is the pixel center in framebuffer pixels. + // Locate the owning grid cell relative to the grid origin + // (top-left = grid_padding.w / .x). + ivec2 orig_grid_pos = ivec2( + floor((gl_FragCoord.xy - uniforms.grid_padding.wx) / uniforms.cell_size) + ); + ivec2 grid_pos = orig_grid_pos; + + // Horizontal padding extend / discard. + if (grid_pos.x < 0) { + if ((uniforms.padding_extend & PAD_EXTEND_LEFT) != 0u) { + grid_pos.x = 0; + } else { + out_color = vec4(0.0); + return; + } + } else if (grid_pos.x > int(uniforms.grid_size.x) - 1) { + if ((uniforms.padding_extend & PAD_EXTEND_RIGHT) != 0u) { + grid_pos.x = int(uniforms.grid_size.x) - 1; + } else { + out_color = vec4(0.0); + return; + } + } + + // Vertical padding extend / discard. + if (grid_pos.y < 0) { + if ((uniforms.padding_extend & PAD_EXTEND_UP) != 0u) { + grid_pos.y = 0; + } else { + out_color = vec4(0.0); + return; + } + } else if (grid_pos.y > int(uniforms.grid_size.y) - 1) { + if ((uniforms.padding_extend & PAD_EXTEND_DOWN) != 0u) { + grid_pos.y = int(uniforms.grid_size.y) - 1; + } else { + out_color = vec4(0.0); + return; + } + } + + // Cursor block fill: only when this fragment's *original* grid_pos + // (pre-clamp) matches the cursor cell. Bypassing the clamp here + // keeps the cursor from leaking into the margin on edge rows. + if (uniforms.cursor_bg_color.a > 0.0 + && orig_grid_pos.x == int(uniforms.cursor_pos.x) + && orig_grid_pos.y == int(uniforms.cursor_pos.y)) + { + vec4 c = uniforms.cursor_bg_color; + c.rgb = grid_prepare_output_rgb(c.rgb, uniforms.input_colorspace); + c.rgb *= c.a; + out_color = c; + return; + } + + // Load + decode the CellBg. + uint idx = uint(grid_pos.y) * uniforms.grid_size.x + uint(grid_pos.x); + vec4 color = unpackUnorm4x8(cells[idx]); + color.rgb = grid_prepare_output_rgb(color.rgb, uniforms.input_colorspace); + color.rgb *= color.a; + + out_color = color; +} diff --git a/sugarloaf/src/grid/shaders/grid_bg.vert.glsl b/sugarloaf/src/grid/shaders/grid_bg.vert.glsl new file mode 100644 index 00000000..788827b4 --- /dev/null +++ b/sugarloaf/src/grid/shaders/grid_bg.vert.glsl @@ -0,0 +1,18 @@ +#version 450 + +// Fullscreen triangle, ported from `grid_bg_vertex` in +// `sugarloaf/src/grid/shaders/grid.metal`. One triangle that covers +// the viewport (clipped at the edges); fragment shader does the work. +// +// gl_VertexIndex 0 → (-1, -3) +// gl_VertexIndex 1 → (-1, 1) +// gl_VertexIndex 2 → ( 3, 1) +// +// Drawn with `vkCmdDraw(cmd, 3, 1, 0, 0)` and TRIANGLE_LIST topology. +// No vertex buffers — `gl_VertexIndex` drives everything. + +void main() { + float x = (gl_VertexIndex == 2) ? 3.0 : -1.0; + float y = (gl_VertexIndex == 0) ? -3.0 : 1.0; + gl_Position = vec4(x, y, 0.0, 1.0); +} diff --git a/sugarloaf/src/grid/shaders/grid_text.frag.glsl b/sugarloaf/src/grid/shaders/grid_text.frag.glsl new file mode 100644 index 00000000..72188b62 --- /dev/null +++ b/sugarloaf/src/grid/shaders/grid_text.frag.glsl @@ -0,0 +1,33 @@ +#version 450 + +// Text fragment shader. Two atlases bound — grayscale (R8) for outline +// glyphs, RGBA8 for color emoji. The vertex shader's `out_atlas` flag +// picks which one to read. +// +// Sampling is `texelFetch` (nearest, no filtering) at integer pixel +// coordinates — matches Metal's `coord::pixel` + `filter::nearest`. + +layout(set = 1, binding = 0) uniform sampler2D atlas_grayscale; +layout(set = 1, binding = 1) uniform sampler2D atlas_color; + +layout(location = 0) flat in uint in_atlas; +layout(location = 1) flat in vec4 in_color; +layout(location = 2) in vec2 in_tex_coord; + +layout(location = 0) out vec4 out_color; + +const uint ATLAS_GRAYSCALE = 0u; + +void main() { + ivec2 uv = ivec2(in_tex_coord); + if (in_atlas == ATLAS_GRAYSCALE) { + // Grayscale: sample alpha mask, multiply by per-glyph color. + // Colour is already premultiplied (in_color.rgb *= in_color.a + // in the vertex shader), so the result is also premultiplied. + float a = texelFetch(atlas_grayscale, uv, 0).r; + out_color = in_color * a; + } else { + // Color atlas: sample RGBA premultiplied directly. + out_color = texelFetch(atlas_color, uv, 0); + } +} diff --git a/sugarloaf/src/grid/shaders/grid_text.vert.glsl b/sugarloaf/src/grid/shaders/grid_text.vert.glsl new file mode 100644 index 00000000..b9517000 --- /dev/null +++ b/sugarloaf/src/grid/shaders/grid_text.vert.glsl @@ -0,0 +1,137 @@ +#version 450 + +// Per-instance text (glyph) vertex shader, ported from +// `grid_text_vertex` in `sugarloaf/src/grid/shaders/grid.metal`. +// Each instance is one `CellText` quad (4-vertex triangle strip). +// +// Vertex layout matches `CellText` in `sugarloaf/src/grid/cell.rs` +// (32 bytes, 7 packed attributes). Vulkan attribute formats are +// chosen per the `cell.rs` layout comments: +// loc 0 R32G32_UINT glyph_pos (offset 0) +// loc 1 R32G32_UINT glyph_size (offset 8) +// loc 2 R16G16_SINT bearings (offset 16, sign-ext to ivec2) +// loc 3 R16G16_UINT grid_pos (offset 20, zero-ext to uvec2) +// loc 4 R8G8B8A8_UNORM color (offset 24, → vec4 0..1) +// loc 5 R8_UINT atlas (offset 28) +// loc 6 R8_UINT bools (offset 29) +// +// Triangle-strip vertex order (4 vertices, as `cmd_draw(4, N, ...)`): +// vid 0 → (0, 0) TL +// vid 1 → (1, 0) TR +// vid 2 → (0, 1) BL +// vid 3 → (1, 1) BR +// Same `corner = (vid==1||vid==3, vid==2||vid==3)` trick the Metal +// shader uses. + +layout(set = 0, binding = 0, std140) uniform Uniforms { + mat4 projection; + vec4 grid_padding; + vec4 cursor_color; + vec4 cursor_bg_color; + vec2 cell_size; + uvec2 grid_size; + uvec2 cursor_pos; + uvec2 _pad_cursor; + float min_contrast; + uint flags; + uint padding_extend; + uint input_colorspace; +} uniforms; + +layout(location = 0) in uvec2 in_glyph_pos; +layout(location = 1) in uvec2 in_glyph_size; +layout(location = 2) in ivec2 in_bearings; +layout(location = 3) in uvec2 in_grid_pos; +layout(location = 4) in vec4 in_color; // unorm8 → vec4 0..1 +layout(location = 5) in uint in_atlas; +layout(location = 6) in uint in_bools; + +layout(location = 0) flat out uint out_atlas; +layout(location = 1) flat out vec4 out_color; +layout(location = 2) out vec2 out_tex_coord; + +const uint BOOL_IS_CURSOR_GLYPH = 2u; + +// Colorspace helpers — same as grid_bg.frag.glsl. We need them in the +// vertex stage too so the foreground color goes through the same +// transform as the bg. +vec3 grid_srgb_to_linear(vec3 c) { + vec3 lo = c / 12.92; + vec3 hi = pow((c + 0.055) / 1.055, vec3(2.4)); + return mix(lo, hi, greaterThan(c, vec3(0.04045))); +} +vec3 grid_linear_to_srgb(vec3 c) { + vec3 lo = c * 12.92; + vec3 hi = pow(c, vec3(1.0 / 2.4)) * 1.055 - 0.055; + return mix(lo, hi, greaterThan(c, vec3(0.0031308))); +} +vec3 grid_srgb_to_p3(vec3 linear_srgb) { + return vec3( + dot(linear_srgb, vec3(0.82246197, 0.17753803, 0.0)), + dot(linear_srgb, vec3(0.03319420, 0.96680580, 0.0)), + dot(linear_srgb, vec3(0.01708263, 0.07239744, 0.91051993)) + ); +} +vec3 grid_rec2020_to_p3(vec3 linear_r2020) { + return vec3( + dot(linear_r2020, vec3( 1.34357825, -0.28217967, -0.06139858)), + dot(linear_r2020, vec3(-0.06529745, 1.08782226, -0.02252481)), + dot(linear_r2020, vec3( 0.00282179, -0.02598807, 1.02316628)) + ); +} +vec3 grid_prepare_output_rgb(vec3 srgb, uint cs) { + vec3 lin = grid_srgb_to_linear(srgb); + if (cs == 0u) { + lin = grid_srgb_to_p3(lin); + } else if (cs == 2u) { + lin = grid_rec2020_to_p3(lin); + } + return grid_linear_to_srgb(lin); +} + +void main() { + // Cell origin in pixel space. + vec2 cell_pos = uniforms.cell_size * vec2(in_grid_pos); + + // Quad corner 0..1 from vertex id (4-vertex TRIANGLE_STRIP). + vec2 corner; + corner.x = float(gl_VertexIndex == 1 || gl_VertexIndex == 3); + corner.y = float(gl_VertexIndex == 2 || gl_VertexIndex == 3); + + // Glyph bbox inside the cell. `bearings.y` is from the *bottom* in + // font convention; flip to top-down by subtracting from cell_size.y. + vec2 size = vec2(in_glyph_size); + vec2 offset = vec2(in_bearings); + offset.y = uniforms.cell_size.y - offset.y; + + vec2 quad = cell_pos + size * corner + offset; + quad.x += uniforms.grid_padding.w; // left + quad.y += uniforms.grid_padding.x; // top + + gl_Position = uniforms.projection * vec4(quad, 0.0, 1.0); + + // Atlas tex coord in PIXEL space — we sample with `texelFetch` + // (nearest filter, no normalization needed). + out_tex_coord = vec2(in_glyph_pos) + size * corner; + out_atlas = in_atlas; + + // Foreground color through the same colorspace pipeline as the bg + // pass. `in_color` is already 0..1 from R8G8B8A8_UNORM. + vec4 color = in_color; + color.rgb = grid_prepare_output_rgb(color.rgb, uniforms.input_colorspace); + color.rgb *= color.a; + + // Cursor cell color swap: if this glyph's cell is under the cursor + // and it's *not* the cursor glyph itself, override with cursor_color. + bool is_cursor_pos = + (in_grid_pos.x == uniforms.cursor_pos.x) && + (in_grid_pos.y == uniforms.cursor_pos.y); + if ((in_bools & BOOL_IS_CURSOR_GLYPH) == 0u && is_cursor_pos) { + vec4 c = uniforms.cursor_color; + c.rgb = grid_prepare_output_rgb(c.rgb, uniforms.input_colorspace); + c.rgb *= c.a; + color = c; + } + + out_color = color; +} diff --git a/sugarloaf/src/grid/shaders/ui_text.vert.glsl b/sugarloaf/src/grid/shaders/ui_text.vert.glsl new file mode 100644 index 00000000..ec0c3f7a --- /dev/null +++ b/sugarloaf/src/grid/shaders/ui_text.vert.glsl @@ -0,0 +1,72 @@ +#version 450 + +// UI text vertex shader, ported from `text_vertex` in +// `sugarloaf/src/grid/shaders/grid.metal`. Positions glyphs in +// pixel space (not on a cell grid) — used by overlay code paths +// (tab titles, command palette, search overlay, assistant) via +// `sugarloaf::text::Text`. +// +// Per-instance vertex layout matches `TextInstance` in +// `sugarloaf/src/text.rs` (36 bytes): +// loc 0 R32G32_SFLOAT pos (offset 0) text-box top-left +// loc 1 R32G32_UINT glyph_pos (offset 8) +// loc 2 R32G32_UINT glyph_size (offset 16) +// loc 3 R16G16_SINT bearings (offset 24) +// loc 4 R8G8B8A8_UNORM color (offset 28) +// loc 5 R8_UINT atlas (offset 32) +// +// 4-vertex triangle strip per instance (`vkCmdDraw(4, N, ..)`). +// +// Uniform block: `vec2 viewport` + `vec2 _pad` for std140 16-byte +// alignment. We compute pixel→NDC inline (no projection matrix — +// uniforms are minimal, just the viewport size). + +layout(set = 0, binding = 0, std140) uniform Uniforms { + vec2 viewport; + vec2 _pad; +} uniforms; + +layout(location = 0) in vec2 in_pos; +layout(location = 1) in uvec2 in_glyph_pos; +layout(location = 2) in uvec2 in_glyph_size; +layout(location = 3) in ivec2 in_bearings; +layout(location = 4) in vec4 in_color; // unorm8 → vec4 0..1 +layout(location = 5) in uint in_atlas; + +layout(location = 0) flat out uint out_atlas; +layout(location = 1) flat out vec4 out_color; +layout(location = 2) out vec2 out_tex_coord; + +void main() { + // Quad corner 0..1 from vertex id (4-vertex TRIANGLE_STRIP). + vec2 corner; + corner.x = float(gl_VertexIndex == 1 || gl_VertexIndex == 3); + corner.y = float(gl_VertexIndex == 2 || gl_VertexIndex == 3); + + vec2 size = vec2(in_glyph_size); + vec2 origin = in_pos + vec2(in_bearings); + vec2 quad_px = origin + size * corner; + + // Pixel → NDC (y-up convention). The Vulkan render pass uses a + // negative-height viewport (set in `Sugarloaf::render_vulkan`) + // so the rasterizer flips this back to top-left origin. Same + // formula as the Metal `text_vertex` — see grid.metal. + vec2 ndc = vec2( + (quad_px.x / uniforms.viewport.x) * 2.0 - 1.0, + 1.0 - (quad_px.y / uniforms.viewport.y) * 2.0 + ); + + gl_Position = vec4(ndc, 0.0, 1.0); + + // Atlas tex coord in PIXEL space — fragment shader uses + // `texelFetch` (nearest filter, no normalization needed). + out_tex_coord = vec2(in_glyph_pos) + size * corner; + out_atlas = in_atlas; + + // Premultiplied RGBA. Color path's atlas already returns + // premultiplied bytes; grayscale path's `color * mask_a` in the + // fragment also stays premultiplied. + vec4 color = in_color; + color.rgb *= color.a; + out_color = color; +} diff --git a/sugarloaf/src/grid/vulkan.rs b/sugarloaf/src/grid/vulkan.rs new file mode 100644 index 00000000..313a2fe6 --- /dev/null +++ b/sugarloaf/src/grid/vulkan.rs @@ -0,0 +1,1347 @@ +// Copyright (c) 2023-present, Raphael Amorim. +// +// This source code is licensed under the MIT license found in the +// LICENSE file in the root directory of this source tree. + +//! Native Vulkan backend for the grid renderer. +//! +//! Phase 4: bg + text passes. Mirrors `grid::metal::MetalGridRenderer` +//! in shape so the rioterm emit loop can drive both via the +//! `GridRenderer` enum without per-backend conditionals. +//! +//! Per-frame ring: bg + uniform + fg buffers are sized per +//! `FRAMES_IN_FLIGHT` slot, indexed by `VulkanFrame::slot`. The +//! `acquire_frame` fence wait already proved that slot's GPU work is +//! done, so writing into slot `N`'s buffers from the CPU is safe. +//! +//! Atlas uploads are deferred — `insert_glyph` records pending pixels +//! into a per-atlas queue, and `render` flushes the queue into the +//! frame's command buffer (one staging buffer per slot, copy + +//! barrier, then the text pass reads). Per-glyph synchronous uploads +//! would cost ~1ms/glyph (vkQueueSubmit + fence wait) — way too slow +//! for the first-frame burst of ~ASCII printables. + +use ash::vk; +use rustc_hash::FxHashMap; + +use super::atlas::{AtlasSlot, GlyphKey, RasterizedGlyph}; +use super::cell::{CellBg, CellText, GridUniforms}; +use crate::context::vulkan::{ + allocate_host_visible_buffer_raw, VulkanBuffer, VulkanContext, VulkanImage, + FRAMES_IN_FLIGHT, +}; +use crate::renderer::image_cache::atlas::AtlasAllocator; + +// Compiled at build time by `sugarloaf/build.rs`. Source GLSL lives +// in `sugarloaf/src/grid/shaders/`; edit those, not the .spv. +const BG_VERT_SPV: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/grid_bg.vert.spv")); +const BG_FRAG_SPV: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/grid_bg.frag.spv")); +const TEXT_VERT_SPV: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/grid_text.vert.spv")); +const TEXT_FRAG_SPV: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/grid_text.frag.spv")); + +/// Extra slots appended to `fg_rows` for cursor glyphs. Mirrors the +/// Metal layout so the CPU emit code is byte-identical. +const CURSOR_ROW_SLOTS: usize = 2; + +/// Initial atlas side. 2048² @ R8 = 4 MiB; matches the Metal default. +const ATLAS_SIZE: u16 = 2048; + +// ======================================================================= +// Glyph atlas +// ======================================================================= + +/// One pending glyph upload — `bytes` were copied at insert time, so +/// the rasterizer's buffer can be reused immediately. Drained by +/// `flush_pending_uploads` on the next `render()`. +struct PendingUpload { + x: u16, + y: u16, + w: u16, + h: u16, + bytes: Vec, +} + +/// Glyph atlas: device-local image + slot allocator + key→slot map + +/// pending upload queue + per-slot staging ring. +/// +/// One instance per atlas kind (R8 grayscale, RGBA8 color). Owned by +/// either `VulkanGridRenderer` (per-panel terminal grids) or +/// `sugarloaf::text::Text`'s Vulkan state (UI overlay text); the +/// caller drives uploads via `prepare_uploads(...)` before +/// `cmd_begin_rendering`. +pub struct VulkanGlyphAtlas { + image: VulkanImage, + allocator: AtlasAllocator, + slots: FxHashMap, + bytes_per_pixel: u32, + pending: Vec, + /// True once the image has been transitioned out of `UNDEFINED`. + /// Until the first upload, the texture is still in `UNDEFINED` + /// layout and reading from it would be UB — the descriptor set is + /// bound but the text pipeline only reads when there are + /// instances, and there are no instances until after at least one + /// `insert_glyph + render` cycle. + initialized: bool, + /// Per-slot staging buffer ring. Sized on demand, never shrinks. + /// Reused across frames within a slot — the `acquire_frame` + /// fence wait inside `VulkanContext` proves the previous use of + /// slot N's staging is GPU-complete before the next reuse. + staging: [Option; FRAMES_IN_FLIGHT], + staging_capacity: [usize; FRAMES_IN_FLIGHT], +} + +impl VulkanGlyphAtlas { + pub fn new_grayscale(ctx: &VulkanContext) -> Self { + Self::new(ctx, vk::Format::R8_UNORM, 1) + } + + pub fn new_color(ctx: &VulkanContext) -> Self { + Self::new(ctx, vk::Format::R8G8B8A8_UNORM, 4) + } + + fn new(ctx: &VulkanContext, format: vk::Format, bytes_per_pixel: u32) -> Self { + let image = ctx.allocate_sampled_image( + ATLAS_SIZE as u32, + ATLAS_SIZE as u32, + format, + vk::ImageUsageFlags::TRANSFER_DST | vk::ImageUsageFlags::SAMPLED, + ); + Self { + image, + allocator: AtlasAllocator::new(ATLAS_SIZE, ATLAS_SIZE), + slots: FxHashMap::default(), + bytes_per_pixel, + pending: Vec::new(), + initialized: false, + staging: std::array::from_fn(|_| None), + staging_capacity: [0; FRAMES_IN_FLIGHT], + } + } + + /// Drain `self.pending` into slot `slot`'s staging buffer (growing + /// it if needed), then record `cmd_copy_buffer_to_image` + + /// barriers into `cmd`. Caller MUST be outside a dynamic-rendering + /// pass — Vulkan 1.3 spec + /// `VUID-vkCmdCopyBufferToImage-renderpass` forbids transfer + /// commands inside one. No-op when there are no pending uploads. + /// + /// We take `(device, instance, physical_device)` rather than + /// `&VulkanContext` so the text overlay path can call this + /// without holding an immutable borrow on the context + /// (`Sugarloaf::render_vulkan` keeps `ctx: &mut VulkanContext` + /// for the swapchain acquire/present cycle). + pub fn flush_uploads( + &mut self, + device: &ash::Device, + instance: &ash::Instance, + physical_device: vk::PhysicalDevice, + cmd: vk::CommandBuffer, + slot: usize, + ) { + if self.pending.is_empty() { + return; + } + let total_bytes: usize = self + .pending + .iter() + .map(|p| (p.w as usize) * (p.h as usize) * self.bytes_per_pixel as usize) + .sum(); + + // Grow per-slot staging if needed. The `min(256K)` floor keeps + // us from churning allocations during the first-frame burst. + if total_bytes > self.staging_capacity[slot] { + let new_cap = total_bytes.next_power_of_two().max(256 * 1024); + self.staging[slot] = + Some(crate::context::vulkan::allocate_host_visible_buffer_raw( + device, + instance, + physical_device, + new_cap as u64, + vk::BufferUsageFlags::TRANSFER_SRC, + )); + self.staging_capacity[slot] = new_cap; + } + let staging = self.staging[slot].as_ref().unwrap(); + let staging_ptr = staging.as_mut_ptr(); + let staging_handle = staging.handle(); + + let bpp = self.bytes_per_pixel as usize; + let mut offset: u64 = 0; + let mut copies: Vec = Vec::with_capacity(self.pending.len()); + unsafe { + for upload in self.pending.drain(..) { + let bytes = (upload.w as usize) * (upload.h as usize) * bpp; + std::ptr::copy_nonoverlapping( + upload.bytes.as_ptr(), + staging_ptr.add(offset as usize), + bytes, + ); + copies.push(image_copy_region( + offset, upload.x, upload.y, upload.w, upload.h, + )); + offset += bytes as u64; + } + } + + upload_to_atlas(device, cmd, staging_handle, self, &copies); + } + + #[inline] + pub fn lookup(&self, key: GlyphKey) -> Option { + self.slots.get(&key).copied() + } + + /// Image view bound to this atlas. Used by callers to wire the + /// atlas into their text-pipeline descriptor sets. + #[inline] + pub fn image_view(&self) -> vk::ImageView { + self.image.view() + } + + /// Pack + queue a glyph for upload. Returns `None` when the atlas + /// is full. Bytes are copied into the pending queue, so the + /// caller's `glyph.bytes` slice can be freed/reused immediately. + /// Pixels reach the GPU on the next `render()` flush. + pub fn insert( + &mut self, + key: GlyphKey, + glyph: RasterizedGlyph<'_>, + ) -> Option { + if glyph.width == 0 || glyph.height == 0 { + // Whitespace / control glyphs — record an empty slot so + // lookups don't keep retrying. + let slot = AtlasSlot { + x: 0, + y: 0, + w: 0, + h: 0, + bearing_x: glyph.bearing_x, + bearing_y: glyph.bearing_y, + }; + self.slots.insert(key, slot); + return Some(slot); + } + + let (x, y) = self.allocator.allocate(glyph.width, glyph.height)?; + let slot = AtlasSlot { + x, + y, + w: glyph.width, + h: glyph.height, + bearing_x: glyph.bearing_x, + bearing_y: glyph.bearing_y, + }; + self.slots.insert(key, slot); + self.pending.push(PendingUpload { + x, + y, + w: glyph.width, + h: glyph.height, + bytes: glyph.bytes.to_vec(), + }); + Some(slot) + } + +} + +// ======================================================================= +// Grid renderer +// ======================================================================= + +pub struct VulkanGridRenderer { + device: ash::Device, + /// Cached so `resize` (which only has `&mut self`) can allocate + /// new bg buffers via `allocate_host_visible_buffer_raw` without + /// needing a `&VulkanContext` borrow. + instance: ash::Instance, + physical_device: vk::PhysicalDevice, + + cols: u32, + rows: u32, + + // ---------- bg state ---------- + bg_buffers: [VulkanBuffer; FRAMES_IN_FLIGHT], + bg_dirty: [bool; FRAMES_IN_FLIGHT], + bg_cpu: Vec, + + // ---------- shared uniform state ---------- + uniform_buffers: [VulkanBuffer; FRAMES_IN_FLIGHT], + + // ---------- bg pipeline ---------- + bg_descriptor_pool: vk::DescriptorPool, + bg_descriptor_set_layout: vk::DescriptorSetLayout, + bg_descriptor_sets: [vk::DescriptorSet; FRAMES_IN_FLIGHT], + bg_pipeline_layout: vk::PipelineLayout, + bg_pipeline: vk::Pipeline, + + // ---------- text state ---------- + fg_rows: Vec>, + fg_staging: Vec, + fg_buffers: [Option; FRAMES_IN_FLIGHT], + fg_capacity: [usize; FRAMES_IN_FLIGHT], + fg_live_count: [u32; FRAMES_IN_FLIGHT], + fg_dirty: [bool; FRAMES_IN_FLIGHT], + + // ---------- text pipeline ---------- + text_uniform_descriptor_set_layout: vk::DescriptorSetLayout, + text_atlas_descriptor_set_layout: vk::DescriptorSetLayout, + text_descriptor_pool: vk::DescriptorPool, + text_uniform_descriptor_sets: [vk::DescriptorSet; FRAMES_IN_FLIGHT], + text_atlas_descriptor_set: vk::DescriptorSet, + text_pipeline_layout: vk::PipelineLayout, + text_pipeline: vk::Pipeline, + sampler: vk::Sampler, + + // ---------- atlases ---------- + pub atlas_grayscale: VulkanGlyphAtlas, + pub atlas_color: VulkanGlyphAtlas, + + needs_full_rebuild: bool, +} + +impl VulkanGridRenderer { + pub fn new(ctx: &VulkanContext, cols: u32, rows: u32) -> Self { + let device = ctx.device().clone(); + let instance = ctx.instance().clone(); + let physical_device = ctx.physical_device(); + + // ----- bg + uniforms ----- + let bg_buffers = std::array::from_fn(|_| alloc_bg_buffer(ctx, cols, rows)); + let uniform_buffers = std::array::from_fn(|_| { + ctx.allocate_host_visible_buffer( + std::mem::size_of::() as u64, + vk::BufferUsageFlags::UNIFORM_BUFFER, + ) + }); + + let bg_descriptor_set_layout = create_bg_descriptor_set_layout(&device); + let bg_descriptor_pool = create_bg_descriptor_pool(&device); + let bg_descriptor_sets = allocate_descriptor_sets( + &device, + bg_descriptor_pool, + bg_descriptor_set_layout, + ); + for slot in 0..FRAMES_IN_FLIGHT { + update_bg_descriptor_set( + &device, + bg_descriptor_sets[slot], + &uniform_buffers[slot], + &bg_buffers[slot], + ); + } + let bg_pipeline_layout = + create_pipeline_layout(&device, &[bg_descriptor_set_layout]); + let pipeline_cache = ctx.pipeline_cache(); + let bg_pipeline = create_bg_pipeline( + &device, + pipeline_cache, + bg_pipeline_layout, + ctx.swapchain_format(), + ); + + // ----- text ----- + let atlas_grayscale = VulkanGlyphAtlas::new_grayscale(ctx); + let atlas_color = VulkanGlyphAtlas::new_color(ctx); + let sampler = create_sampler(&device); + + let text_uniform_descriptor_set_layout = + create_text_uniform_descriptor_set_layout(&device); + let text_atlas_descriptor_set_layout = + create_text_atlas_descriptor_set_layout(&device); + // One pool that holds (FRAMES_IN_FLIGHT uniform sets) + (1 atlas set). + let text_descriptor_pool = create_text_descriptor_pool(&device); + + let text_uniform_descriptor_sets = allocate_descriptor_sets( + &device, + text_descriptor_pool, + text_uniform_descriptor_set_layout, + ); + for slot in 0..FRAMES_IN_FLIGHT { + update_text_uniform_descriptor_set( + &device, + text_uniform_descriptor_sets[slot], + &uniform_buffers[slot], + ); + } + let text_atlas_descriptor_set = allocate_one_descriptor_set( + &device, + text_descriptor_pool, + text_atlas_descriptor_set_layout, + ); + update_text_atlas_descriptor_set( + &device, + text_atlas_descriptor_set, + &atlas_grayscale.image, + &atlas_color.image, + sampler, + ); + + let text_pipeline_layout = create_pipeline_layout( + &device, + &[ + text_uniform_descriptor_set_layout, + text_atlas_descriptor_set_layout, + ], + ); + let text_pipeline = create_text_pipeline( + &device, + pipeline_cache, + text_pipeline_layout, + ctx.swapchain_format(), + ); + + let bg_len = (cols as usize) * (rows as usize); + Self { + device, + instance, + physical_device, + cols, + rows, + bg_buffers, + bg_dirty: [true; FRAMES_IN_FLIGHT], + bg_cpu: vec![CellBg::TRANSPARENT; bg_len], + uniform_buffers, + bg_descriptor_pool, + bg_descriptor_set_layout, + bg_descriptor_sets, + bg_pipeline_layout, + bg_pipeline, + fg_rows: init_fg_rows(rows), + fg_staging: Vec::new(), + fg_buffers: std::array::from_fn(|_| None), + fg_capacity: [0; FRAMES_IN_FLIGHT], + fg_live_count: [0; FRAMES_IN_FLIGHT], + fg_dirty: [true; FRAMES_IN_FLIGHT], + text_uniform_descriptor_set_layout, + text_atlas_descriptor_set_layout, + text_descriptor_pool, + text_uniform_descriptor_sets, + text_atlas_descriptor_set, + text_pipeline_layout, + text_pipeline, + sampler, + atlas_grayscale, + atlas_color, + needs_full_rebuild: true, + } + } + + #[inline] + pub fn needs_full_rebuild(&self) -> bool { + self.needs_full_rebuild + } + + #[inline] + pub fn mark_full_rebuild_done(&mut self) { + self.needs_full_rebuild = false; + } + + pub fn resize(&mut self, cols: u32, rows: u32) { + if cols == self.cols && rows == self.rows { + return; + } + unsafe { + let _ = self.device.device_wait_idle(); + } + + self.cols = cols; + self.rows = rows; + let bg_len = (cols as usize) * (rows as usize); + self.bg_cpu = vec![CellBg::TRANSPARENT; bg_len]; + + // Reallocate bg buffers via the cached (instance, + // physical_device) pair and re-wire descriptor sets to the + // new buffer handles. + let bg_byte_size = (bg_len * std::mem::size_of::()) + .max(std::mem::size_of::()) as u64; + self.bg_buffers = std::array::from_fn(|_| { + allocate_host_visible_buffer_raw( + &self.device, + &self.instance, + self.physical_device, + bg_byte_size, + vk::BufferUsageFlags::STORAGE_BUFFER, + ) + }); + for slot in 0..FRAMES_IN_FLIGHT { + update_bg_descriptor_set( + &self.device, + self.bg_descriptor_sets[slot], + &self.uniform_buffers[slot], + &self.bg_buffers[slot], + ); + } + self.bg_dirty = [true; FRAMES_IN_FLIGHT]; + + // Reset fg state — emit loop will re-populate after resize. + self.fg_rows = init_fg_rows(rows); + self.fg_dirty = [true; FRAMES_IN_FLIGHT]; + self.fg_live_count = [0; FRAMES_IN_FLIGHT]; + self.needs_full_rebuild = true; + } + + pub fn write_row(&mut self, row: u32, bg: &[CellBg], fg: &[CellText]) { + // FG: stash in CPU per-row vec, mark all slots dirty. + 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; FRAMES_IN_FLIGHT]; + } + + if row >= self.rows { + return; + } + let row_start = (row as usize) * (self.cols as usize); + let row_len = (self.cols as usize).min(bg.len()); + self.bg_cpu[row_start..row_start + row_len].copy_from_slice(&bg[..row_len]); + for slot in &mut self.bg_cpu[row_start + row_len..row_start + self.cols as usize] + { + *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; FRAMES_IN_FLIGHT]; + } + slot.clear(); + } + if row >= self.rows { + return; + } + let row_start = (row as usize) * (self.cols as usize); + for slot in &mut self.bg_cpu[row_start..row_start + self.cols as usize] { + *slot = CellBg::TRANSPARENT; + } + self.bg_dirty = [true; FRAMES_IN_FLIGHT]; + } + + #[inline] + pub fn lookup_glyph(&self, key: GlyphKey) -> Option { + self.atlas_grayscale.lookup(key) + } + + #[inline] + pub fn lookup_glyph_color(&self, key: GlyphKey) -> Option { + self.atlas_color.lookup(key) + } + + #[inline] + pub fn insert_glyph( + &mut self, + key: GlyphKey, + glyph: RasterizedGlyph<'_>, + ) -> Option { + self.atlas_grayscale.insert(key, glyph) + } + + #[inline] + pub fn insert_glyph_color( + &mut self, + key: GlyphKey, + glyph: RasterizedGlyph<'_>, + ) -> Option { + self.atlas_color.insert(key, glyph) + } + + /// Drain pending atlas uploads into `cmd`. MUST be called BEFORE + /// `Sugarloaf::render_vulkan` opens its dynamic-rendering pass — + /// `vkCmdCopyBufferToImage` is forbidden inside a render pass. + /// No-op when both atlases have no pending entries. + pub fn prepare( + &mut self, + ctx: &VulkanContext, + cmd: vk::CommandBuffer, + frame_slot: usize, + ) { + debug_assert!(frame_slot < FRAMES_IN_FLIGHT); + if self.atlas_grayscale.pending.is_empty() && self.atlas_color.pending.is_empty() + { + return; + } + 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( + &mut self, + ctx: &VulkanContext, + cmd: vk::CommandBuffer, + frame_slot: usize, + uniforms: &GridUniforms, + ) { + 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; + std::ptr::copy_nonoverlapping( + self.bg_cpu.as_ptr(), + dst, + self.bg_cpu.len(), + ); + } + self.bg_dirty[slot] = false; + } + unsafe { + let dst = self.uniform_buffers[slot].as_mut_ptr() as *mut GridUniforms; + std::ptr::write(dst, *uniforms); + } + + // ----- bg pass (1 fullscreen triangle, fragment does cell lookup) ----- + unsafe { + self.device.cmd_bind_pipeline( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.bg_pipeline, + ); + self.device.cmd_bind_descriptor_sets( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.bg_pipeline_layout, + 0, + &[self.bg_descriptor_sets[slot]], + &[], + ); + self.device.cmd_draw(cmd, 3, 1, 0, 0); + } + + // ----- text pass (instanced quads, one per glyph) ----- + if self.fg_dirty[slot] { + self.fg_staging.clear(); + for row in &self.fg_rows { + self.fg_staging.extend_from_slice(row); + } + let needed = self.fg_staging.len(); + + if needed > self.fg_capacity[slot] { + let new_cap = needed.next_power_of_two().max(64); + self.fg_buffers[slot] = Some(ctx.allocate_host_visible_buffer( + (new_cap * std::mem::size_of::()) as u64, + vk::BufferUsageFlags::VERTEX_BUFFER, + )); + self.fg_capacity[slot] = new_cap; + } + + if needed > 0 { + let buf = self.fg_buffers[slot].as_ref().unwrap(); + unsafe { + let dst = buf.as_mut_ptr() as *mut CellText; + std::ptr::copy_nonoverlapping(self.fg_staging.as_ptr(), dst, needed); + } + } + self.fg_live_count[slot] = needed as u32; + self.fg_dirty[slot] = false; + } + + let instance_count = self.fg_live_count[slot]; + if instance_count == 0 { + return; + } + + unsafe { + self.device.cmd_bind_pipeline( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.text_pipeline, + ); + self.device.cmd_bind_descriptor_sets( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.text_pipeline_layout, + 0, + &[ + self.text_uniform_descriptor_sets[slot], + self.text_atlas_descriptor_set, + ], + &[], + ); + let buf = self.fg_buffers[slot].as_ref().unwrap(); + self.device + .cmd_bind_vertex_buffers(cmd, 0, &[buf.handle()], &[0]); + self.device.cmd_draw(cmd, 4, instance_count, 0, 0); + } + } + + /// Delegate to each atlas's own `flush_uploads`. Each atlas owns + /// its own per-slot staging buffer ring now — see + /// `VulkanGlyphAtlas::flush_uploads`. + fn flush_pending_uploads( + &mut self, + _ctx: &VulkanContext, + cmd: vk::CommandBuffer, + slot: usize, + ) { + self.atlas_grayscale.flush_uploads( + &self.device, + &self.instance, + self.physical_device, + cmd, + slot, + ); + self.atlas_color.flush_uploads( + &self.device, + &self.instance, + self.physical_device, + cmd, + slot, + ); + } +} + +/// Record an atlas upload: barrier image → `TRANSFER_DST_OPTIMAL`, +/// `cmd_copy_buffer_to_image`, barrier image → `SHADER_READ_ONLY_OPTIMAL`. +/// +/// Both barriers are required: the first synchronizes any prior +/// fragment-shader read of the atlas (steady state) against the +/// upcoming transfer write; the second synchronizes the transfer +/// write against the *next* fragment-shader read (which happens in +/// the same command buffer, in the text pipeline draw a few hundred +/// instructions later). Without the trailing barrier the GPU is free +/// to start the fragment work before the copy completes, producing +/// transient garbage glyphs. +/// +/// Caller (`flush_pending_uploads`) must ensure this is invoked +/// *outside* a dynamic-rendering pass — Vulkan 1.3 spec +/// VUID-vkCmdCopyBufferToImage-renderpass forbids transfer commands +/// inside a render pass. `Sugarloaf::render_vulkan` honours this by +/// calling `prepare_vulkan` before `cmd_begin_rendering`. +fn upload_to_atlas( + device: &ash::Device, + cmd: vk::CommandBuffer, + staging: vk::Buffer, + atlas: &mut VulkanGlyphAtlas, + copies: &[vk::BufferImageCopy], +) { + let old_layout = if atlas.initialized { + vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL + } else { + vk::ImageLayout::UNDEFINED + }; + unsafe { + // → TRANSFER_DST + let to_transfer = vk::ImageMemoryBarrier2::default() + .src_stage_mask(if atlas.initialized { + vk::PipelineStageFlags2::FRAGMENT_SHADER + } else { + vk::PipelineStageFlags2::TOP_OF_PIPE + }) + .src_access_mask(if atlas.initialized { + vk::AccessFlags2::SHADER_READ + } else { + vk::AccessFlags2::empty() + }) + .dst_stage_mask(vk::PipelineStageFlags2::COPY) + .dst_access_mask(vk::AccessFlags2::TRANSFER_WRITE) + .old_layout(old_layout) + .new_layout(vk::ImageLayout::TRANSFER_DST_OPTIMAL) + .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .image(atlas.image.handle()) + .subresource_range(color_subresource_range()); + let barriers = [to_transfer]; + let dep = vk::DependencyInfo::default().image_memory_barriers(&barriers); + device.cmd_pipeline_barrier2(cmd, &dep); + + // copy + device.cmd_copy_buffer_to_image( + cmd, + staging, + atlas.image.handle(), + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + copies, + ); + + // → SHADER_READ + let to_shader_read = vk::ImageMemoryBarrier2::default() + .src_stage_mask(vk::PipelineStageFlags2::COPY) + .src_access_mask(vk::AccessFlags2::TRANSFER_WRITE) + .dst_stage_mask(vk::PipelineStageFlags2::FRAGMENT_SHADER) + .dst_access_mask(vk::AccessFlags2::SHADER_READ) + .old_layout(vk::ImageLayout::TRANSFER_DST_OPTIMAL) + .new_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL) + .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .image(atlas.image.handle()) + .subresource_range(color_subresource_range()); + let barriers = [to_shader_read]; + let dep = vk::DependencyInfo::default().image_memory_barriers(&barriers); + device.cmd_pipeline_barrier2(cmd, &dep); + } + + atlas.initialized = true; +} + +impl Drop for VulkanGridRenderer { + fn drop(&mut self) { + unsafe { + let _ = self.device.device_wait_idle(); + self.device.destroy_pipeline(self.text_pipeline, None); + self.device + .destroy_pipeline_layout(self.text_pipeline_layout, None); + self.device + .destroy_descriptor_pool(self.text_descriptor_pool, None); + self.device.destroy_descriptor_set_layout( + self.text_atlas_descriptor_set_layout, + None, + ); + self.device.destroy_descriptor_set_layout( + self.text_uniform_descriptor_set_layout, + None, + ); + self.device.destroy_sampler(self.sampler, None); + + self.device.destroy_pipeline(self.bg_pipeline, None); + self.device + .destroy_pipeline_layout(self.bg_pipeline_layout, None); + self.device + .destroy_descriptor_pool(self.bg_descriptor_pool, None); + self.device + .destroy_descriptor_set_layout(self.bg_descriptor_set_layout, None); + // Buffers + atlas images drop themselves. + } + } +} + +// ======================================================================= +// Helpers +// ======================================================================= + +fn alloc_bg_buffer(ctx: &VulkanContext, cols: u32, rows: u32) -> VulkanBuffer { + let size = (cols as u64) + .saturating_mul(rows as u64) + .saturating_mul(std::mem::size_of::() as u64) + .max(std::mem::size_of::() as u64); + ctx.allocate_host_visible_buffer(size, vk::BufferUsageFlags::STORAGE_BUFFER) +} + +fn init_fg_rows(rows: u32) -> Vec> { + (0..(rows as usize + CURSOR_ROW_SLOTS)) + .map(|_| Vec::new()) + .collect() +} + +fn color_subresource_range() -> vk::ImageSubresourceRange { + vk::ImageSubresourceRange::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .base_mip_level(0) + .level_count(1) + .base_array_layer(0) + .layer_count(1) +} + +fn image_copy_region( + buffer_offset: u64, + x: u16, + y: u16, + w: u16, + h: u16, +) -> vk::BufferImageCopy { + vk::BufferImageCopy::default() + .buffer_offset(buffer_offset) + .buffer_row_length(0) // tightly packed — same as bytes_per_row = w * bpp + .buffer_image_height(0) // tightly packed + .image_subresource( + vk::ImageSubresourceLayers::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .mip_level(0) + .base_array_layer(0) + .layer_count(1), + ) + .image_offset(vk::Offset3D { + x: x as i32, + y: y as i32, + z: 0, + }) + .image_extent(vk::Extent3D { + width: w as u32, + height: h as u32, + depth: 1, + }) +} + +// ----- descriptor / pipeline setup helpers ----- + +fn create_bg_descriptor_set_layout(device: &ash::Device) -> vk::DescriptorSetLayout { + let bindings = [ + vk::DescriptorSetLayoutBinding::default() + .binding(0) + .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::VERTEX | vk::ShaderStageFlags::FRAGMENT), + vk::DescriptorSetLayoutBinding::default() + .binding(1) + .descriptor_type(vk::DescriptorType::STORAGE_BUFFER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::FRAGMENT), + ]; + let info = vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings); + unsafe { + device + .create_descriptor_set_layout(&info, None) + .expect("create_descriptor_set_layout(grid.bg)") + } +} + +fn create_bg_descriptor_pool(device: &ash::Device) -> vk::DescriptorPool { + let sizes = [ + vk::DescriptorPoolSize { + ty: vk::DescriptorType::UNIFORM_BUFFER, + descriptor_count: FRAMES_IN_FLIGHT as u32, + }, + vk::DescriptorPoolSize { + ty: vk::DescriptorType::STORAGE_BUFFER, + descriptor_count: FRAMES_IN_FLIGHT as u32, + }, + ]; + let info = vk::DescriptorPoolCreateInfo::default() + .max_sets(FRAMES_IN_FLIGHT as u32) + .pool_sizes(&sizes); + unsafe { + device + .create_descriptor_pool(&info, None) + .expect("create_descriptor_pool(grid.bg)") + } +} + +fn create_text_uniform_descriptor_set_layout( + device: &ash::Device, +) -> vk::DescriptorSetLayout { + let bindings = [vk::DescriptorSetLayoutBinding::default() + .binding(0) + .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::VERTEX | vk::ShaderStageFlags::FRAGMENT)]; + let info = vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings); + unsafe { + device + .create_descriptor_set_layout(&info, None) + .expect("create_descriptor_set_layout(grid.text uniform)") + } +} + +fn create_text_atlas_descriptor_set_layout( + device: &ash::Device, +) -> vk::DescriptorSetLayout { + let bindings = [ + vk::DescriptorSetLayoutBinding::default() + .binding(0) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::FRAGMENT), + vk::DescriptorSetLayoutBinding::default() + .binding(1) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::FRAGMENT), + ]; + let info = vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings); + unsafe { + device + .create_descriptor_set_layout(&info, None) + .expect("create_descriptor_set_layout(grid.text atlas)") + } +} + +fn create_text_descriptor_pool(device: &ash::Device) -> vk::DescriptorPool { + let sizes = [ + vk::DescriptorPoolSize { + ty: vk::DescriptorType::UNIFORM_BUFFER, + descriptor_count: FRAMES_IN_FLIGHT as u32, + }, + vk::DescriptorPoolSize { + ty: vk::DescriptorType::COMBINED_IMAGE_SAMPLER, + descriptor_count: 2, + }, + ]; + let info = vk::DescriptorPoolCreateInfo::default() + .max_sets((FRAMES_IN_FLIGHT + 1) as u32) + .pool_sizes(&sizes); + unsafe { + device + .create_descriptor_pool(&info, None) + .expect("create_descriptor_pool(grid.text)") + } +} + +fn allocate_descriptor_sets( + device: &ash::Device, + pool: vk::DescriptorPool, + layout: vk::DescriptorSetLayout, +) -> [vk::DescriptorSet; FRAMES_IN_FLIGHT] { + let layouts = [layout; FRAMES_IN_FLIGHT]; + let info = vk::DescriptorSetAllocateInfo::default() + .descriptor_pool(pool) + .set_layouts(&layouts); + let sets = unsafe { + device + .allocate_descriptor_sets(&info) + .expect("allocate_descriptor_sets") + }; + let mut out = [vk::DescriptorSet::null(); FRAMES_IN_FLIGHT]; + out.copy_from_slice(&sets); + out +} + +fn allocate_one_descriptor_set( + device: &ash::Device, + pool: vk::DescriptorPool, + layout: vk::DescriptorSetLayout, +) -> vk::DescriptorSet { + let layouts = [layout]; + let info = vk::DescriptorSetAllocateInfo::default() + .descriptor_pool(pool) + .set_layouts(&layouts); + unsafe { + device + .allocate_descriptor_sets(&info) + .expect("allocate_descriptor_sets(one)")[0] + } +} + +fn update_bg_descriptor_set( + device: &ash::Device, + set: vk::DescriptorSet, + uniform: &VulkanBuffer, + cells: &VulkanBuffer, +) { + let uniform_info = vk::DescriptorBufferInfo::default() + .buffer(uniform.handle()) + .offset(0) + .range(uniform.size()); + let uniform_infos = [uniform_info]; + let cells_info = vk::DescriptorBufferInfo::default() + .buffer(cells.handle()) + .offset(0) + .range(cells.size()); + let cells_infos = [cells_info]; + + let writes = [ + vk::WriteDescriptorSet::default() + .dst_set(set) + .dst_binding(0) + .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER) + .buffer_info(&uniform_infos), + vk::WriteDescriptorSet::default() + .dst_set(set) + .dst_binding(1) + .descriptor_type(vk::DescriptorType::STORAGE_BUFFER) + .buffer_info(&cells_infos), + ]; + unsafe { + device.update_descriptor_sets(&writes, &[]); + } +} + +fn update_text_uniform_descriptor_set( + device: &ash::Device, + set: vk::DescriptorSet, + uniform: &VulkanBuffer, +) { + let uniform_info = vk::DescriptorBufferInfo::default() + .buffer(uniform.handle()) + .offset(0) + .range(uniform.size()); + let uniform_infos = [uniform_info]; + let writes = [vk::WriteDescriptorSet::default() + .dst_set(set) + .dst_binding(0) + .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER) + .buffer_info(&uniform_infos)]; + unsafe { + device.update_descriptor_sets(&writes, &[]); + } +} + +fn update_text_atlas_descriptor_set( + device: &ash::Device, + set: vk::DescriptorSet, + grayscale: &VulkanImage, + color: &VulkanImage, + sampler: vk::Sampler, +) { + let gray_info = vk::DescriptorImageInfo::default() + .sampler(sampler) + .image_view(grayscale.view()) + .image_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL); + let gray_infos = [gray_info]; + let color_info = vk::DescriptorImageInfo::default() + .sampler(sampler) + .image_view(color.view()) + .image_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL); + let color_infos = [color_info]; + let writes = [ + vk::WriteDescriptorSet::default() + .dst_set(set) + .dst_binding(0) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .image_info(&gray_infos), + vk::WriteDescriptorSet::default() + .dst_set(set) + .dst_binding(1) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .image_info(&color_infos), + ]; + unsafe { + device.update_descriptor_sets(&writes, &[]); + } +} + +fn create_pipeline_layout( + device: &ash::Device, + set_layouts: &[vk::DescriptorSetLayout], +) -> vk::PipelineLayout { + let info = vk::PipelineLayoutCreateInfo::default().set_layouts(set_layouts); + unsafe { + device + .create_pipeline_layout(&info, None) + .expect("create_pipeline_layout(grid)") + } +} + +fn create_sampler(device: &ash::Device) -> vk::Sampler { + // Nearest filter + clamp-to-edge — matches Metal's + // `filter::nearest, address::clamp_to_edge`. Not used for + // sampling per se (we use `texelFetch` in the fragment shader), + // but the COMBINED_IMAGE_SAMPLER descriptor still requires a + // sampler object. + let info = vk::SamplerCreateInfo::default() + .mag_filter(vk::Filter::NEAREST) + .min_filter(vk::Filter::NEAREST) + .mipmap_mode(vk::SamplerMipmapMode::NEAREST) + .address_mode_u(vk::SamplerAddressMode::CLAMP_TO_EDGE) + .address_mode_v(vk::SamplerAddressMode::CLAMP_TO_EDGE) + .address_mode_w(vk::SamplerAddressMode::CLAMP_TO_EDGE); + unsafe { + device + .create_sampler(&info, None) + .expect("create_sampler(grid.text)") + } +} + +fn create_bg_pipeline( + device: &ash::Device, + pipeline_cache: vk::PipelineCache, + layout: vk::PipelineLayout, + color_format: vk::Format, +) -> vk::Pipeline { + build_pipeline( + device, + pipeline_cache, + layout, + color_format, + BG_VERT_SPV, + BG_FRAG_SPV, + &[], // no vertex bindings + &[], + vk::PrimitiveTopology::TRIANGLE_LIST, + BlendMode::Premultiplied, // bg uses src=SRC_ALPHA + ) +} + +fn create_text_pipeline( + device: &ash::Device, + pipeline_cache: vk::PipelineCache, + layout: vk::PipelineLayout, + color_format: vk::Format, +) -> vk::Pipeline { + let bindings = [vk::VertexInputBindingDescription::default() + .binding(0) + .stride(std::mem::size_of::() as u32) + .input_rate(vk::VertexInputRate::INSTANCE)]; + let attrs = [ + // 0: glyph_pos uvec2 @ 0 + vk::VertexInputAttributeDescription::default() + .location(0) + .binding(0) + .format(vk::Format::R32G32_UINT) + .offset(0), + // 1: glyph_size uvec2 @ 8 + vk::VertexInputAttributeDescription::default() + .location(1) + .binding(0) + .format(vk::Format::R32G32_UINT) + .offset(8), + // 2: bearings ivec2 @ 16 (stored as i16x2) + vk::VertexInputAttributeDescription::default() + .location(2) + .binding(0) + .format(vk::Format::R16G16_SINT) + .offset(16), + // 3: grid_pos uvec2 @ 20 (stored as u16x2) + vk::VertexInputAttributeDescription::default() + .location(3) + .binding(0) + .format(vk::Format::R16G16_UINT) + .offset(20), + // 4: color vec4 @ 24 (UNORM8) + vk::VertexInputAttributeDescription::default() + .location(4) + .binding(0) + .format(vk::Format::R8G8B8A8_UNORM) + .offset(24), + // 5: atlas u8 @ 28 → uint + vk::VertexInputAttributeDescription::default() + .location(5) + .binding(0) + .format(vk::Format::R8_UINT) + .offset(28), + // 6: bools u8 @ 29 → uint + vk::VertexInputAttributeDescription::default() + .location(6) + .binding(0) + .format(vk::Format::R8_UINT) + .offset(29), + ]; + build_pipeline( + device, + pipeline_cache, + layout, + color_format, + TEXT_VERT_SPV, + TEXT_FRAG_SPV, + &bindings, + &attrs, + vk::PrimitiveTopology::TRIANGLE_STRIP, + BlendMode::PremultipliedOverFromOne, // text fragment returns premultiplied + ) +} + +#[derive(Copy, Clone)] +enum BlendMode { + /// Source RGB factor = `SRC_ALPHA`. For shaders that return + /// non-premultiplied RGBA + alpha (the bg pass). + Premultiplied, + /// Source RGB factor = `ONE`. For shaders that return + /// already-premultiplied RGBA (the text pass — `in.color * mask_a` + /// and the color atlas sample are both premultiplied). + PremultipliedOverFromOne, +} + +#[allow(clippy::too_many_arguments)] +fn build_pipeline( + device: &ash::Device, + pipeline_cache: vk::PipelineCache, + layout: vk::PipelineLayout, + color_format: vk::Format, + vert_spv: &[u8], + frag_spv: &[u8], + vertex_bindings: &[vk::VertexInputBindingDescription], + vertex_attrs: &[vk::VertexInputAttributeDescription], + topology: vk::PrimitiveTopology, + blend: BlendMode, +) -> vk::Pipeline { + let vert = load_shader_module(device, vert_spv); + let frag = load_shader_module(device, frag_spv); + + let entry = c"main"; + let stages = [ + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::VERTEX) + .module(vert) + .name(entry), + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::FRAGMENT) + .module(frag) + .name(entry), + ]; + + let vertex_input = vk::PipelineVertexInputStateCreateInfo::default() + .vertex_binding_descriptions(vertex_bindings) + .vertex_attribute_descriptions(vertex_attrs); + + let input_assembly = vk::PipelineInputAssemblyStateCreateInfo::default() + .topology(topology) + .primitive_restart_enable(false); + + let viewport_state = vk::PipelineViewportStateCreateInfo::default() + .viewport_count(1) + .scissor_count(1); + + let rasterization = vk::PipelineRasterizationStateCreateInfo::default() + .polygon_mode(vk::PolygonMode::FILL) + .cull_mode(vk::CullModeFlags::NONE) + .front_face(vk::FrontFace::COUNTER_CLOCKWISE) + .line_width(1.0); + + let multisample = vk::PipelineMultisampleStateCreateInfo::default() + .rasterization_samples(vk::SampleCountFlags::TYPE_1); + + let (src_rgb, dst_rgb) = match blend { + BlendMode::Premultiplied => ( + vk::BlendFactor::SRC_ALPHA, + vk::BlendFactor::ONE_MINUS_SRC_ALPHA, + ), + BlendMode::PremultipliedOverFromOne => { + (vk::BlendFactor::ONE, vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + } + }; + let blend_attachment = vk::PipelineColorBlendAttachmentState::default() + .blend_enable(true) + .src_color_blend_factor(src_rgb) + .dst_color_blend_factor(dst_rgb) + .color_blend_op(vk::BlendOp::ADD) + .src_alpha_blend_factor(vk::BlendFactor::ONE) + .dst_alpha_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .alpha_blend_op(vk::BlendOp::ADD) + .color_write_mask(vk::ColorComponentFlags::RGBA); + let blend_attachments = [blend_attachment]; + let color_blend = + vk::PipelineColorBlendStateCreateInfo::default().attachments(&blend_attachments); + + let dynamic_states = [vk::DynamicState::VIEWPORT, vk::DynamicState::SCISSOR]; + let dynamic_state = + vk::PipelineDynamicStateCreateInfo::default().dynamic_states(&dynamic_states); + + let color_attachment_formats = [color_format]; + let mut rendering = vk::PipelineRenderingCreateInfo::default() + .color_attachment_formats(&color_attachment_formats); + + let pipeline_info = vk::GraphicsPipelineCreateInfo::default() + .stages(&stages) + .vertex_input_state(&vertex_input) + .input_assembly_state(&input_assembly) + .viewport_state(&viewport_state) + .rasterization_state(&rasterization) + .multisample_state(&multisample) + .color_blend_state(&color_blend) + .dynamic_state(&dynamic_state) + .layout(layout) + .push_next(&mut rendering); + + let pipeline = unsafe { + device + .create_graphics_pipelines(pipeline_cache, &[pipeline_info], None) + .map_err(|(_, e)| e) + .expect("create_graphics_pipelines(grid)")[0] + }; + + unsafe { + device.destroy_shader_module(vert, None); + device.destroy_shader_module(frag, None); + } + pipeline +} + +fn load_shader_module(device: &ash::Device, bytes: &[u8]) -> vk::ShaderModule { + let code = ash::util::read_spv(&mut std::io::Cursor::new(bytes)) + .expect("read_spv (embedded grid shader is valid)"); + let info = vk::ShaderModuleCreateInfo::default().code(&code); + unsafe { + device + .create_shader_module(&info, None) + .expect("create_shader_module(grid)") + } +} diff --git a/sugarloaf/src/lib.rs b/sugarloaf/src/lib.rs index 65f083bc..07c4c4a4 100644 --- a/sugarloaf/src/lib.rs +++ b/sugarloaf/src/lib.rs @@ -12,7 +12,9 @@ pub mod text; // This path was used by the in-tree fork; preserve it for stability. pub use swash; -// Expose WGPU +// Expose WGPU when the `wgpu` feature is enabled. Downstream code +// that needs `wgpu::Color` etc. picks it up via `sugarloaf::wgpu::…`. +#[cfg(feature = "wgpu")] pub use wgpu; pub use swash::{Attributes, Stretch, Style, Weight}; @@ -28,9 +30,11 @@ pub use crate::sugarloaf::{ CursorKind, DrawableChar, ImageProperties, Object, Quad, Rect, RichText, RichTextLinesRange, RichTextRenderData, SugarCursor, }, - Colorspace, Sugarloaf, SugarloafBackend, SugarloafErrors, SugarloafRenderer, + Color, Colorspace, Sugarloaf, SugarloafBackend, SugarloafErrors, SugarloafRenderer, SugarloafWindow, SugarloafWindowSize, SugarloafWithErrors, }; +// `Filter` is the librashader CRT/scanline-filter wrapper — wgpu-only. +#[cfg(feature = "wgpu")] pub use components::filters::Filter; pub use layout::{ Content, RichTextConfig, SpanStyle, SpanStyleDecoration, TextDimensions, diff --git a/sugarloaf/src/renderer/clear.frag.glsl b/sugarloaf/src/renderer/clear.frag.glsl new file mode 100644 index 00000000..0d9450a5 --- /dev/null +++ b/sugarloaf/src/renderer/clear.frag.glsl @@ -0,0 +1,11 @@ +#version 450 + +layout(push_constant) uniform PC { + vec4 color; +} pc; + +layout(location = 0) out vec4 out_color; + +void main() { + out_color = pc.color; +} diff --git a/sugarloaf/src/renderer/clear.vert.glsl b/sugarloaf/src/renderer/clear.vert.glsl new file mode 100644 index 00000000..2ca8aed0 --- /dev/null +++ b/sugarloaf/src/renderer/clear.vert.glsl @@ -0,0 +1,15 @@ +#version 450 + +// Bootstrap pipeline: generates a centered rect via gl_VertexIndex with +// TRIANGLE_STRIP topology. No vertex buffers; no descriptor sets. This +// shader is temporary scaffolding to prove the Vulkan pipeline plumbing +// works end-to-end — the real renderer.vert.glsl ports `renderer.metal`. +void main() { + // index bit 0 -> x in {0, 1}; bit 1 -> y in {0, 1} + // 0: (-0.5, -0.5) 2: (-0.5, 0.5) + // 1: ( 0.5, -0.5) 3: ( 0.5, 0.5) + // Strip order (0,1,2,3) -> two triangles forming a centered rect. + float x = float(gl_VertexIndex & 1) - 0.5; + float y = float((gl_VertexIndex >> 1) & 1) - 0.5; + gl_Position = vec4(x, y, 0.0, 1.0); +} diff --git a/sugarloaf/src/renderer/cpu.rs b/sugarloaf/src/renderer/cpu.rs index 5b7f6b34..1d52e9a6 100644 --- a/sugarloaf/src/renderer/cpu.rs +++ b/sugarloaf/src/renderer/cpu.rs @@ -290,7 +290,7 @@ pub fn render_cpu( ctx: &mut CpuContext, renderer: &Renderer, cache: &mut CpuCache, - background: Option, + background: Option, ) { let vertices = renderer.vertices(); diff --git a/sugarloaf/src/renderer/geometry.vert.glsl b/sugarloaf/src/renderer/geometry.vert.glsl new file mode 100644 index 00000000..0e3d92ca --- /dev/null +++ b/sugarloaf/src/renderer/geometry.vert.glsl @@ -0,0 +1,57 @@ +#version 450 + +// Per-vertex non-quad geometry shader, ported from `vs_main` in +// `sugarloaf/src/renderer/renderer.metal`. Used for `polygon()` / +// `line()` / `triangle()` / `arc()` calls — anything that emits +// caller-supplied vertices instead of QuadInstances. +// +// Per-vertex layout matches `Vertex` in +// `sugarloaf/src/renderer/batch.rs` (88 bytes): +// loc 0 R32G32B32_SFLOAT pos (offset 0) +// loc 1 R32G32B32A32_SFLOAT color (offset 12) +// loc 2 R32G32_SFLOAT uv (offset 28) +// loc 3 R32G32_SINT layers (offset 36) +// loc 4 R32G32B32A32_SFLOAT corner_radii (offset 44) +// loc 5 R32G32_SFLOAT rect_size (offset 60) +// loc 6 R32_SINT underline_style (offset 68) +// loc 7 R32G32B32A32_SFLOAT clip_rect (offset 72) +// +// Drawn as TRIANGLE_LIST with one instance — the caller supplies +// every vertex explicitly. Outputs the same `VertexOutput` as +// `quad.vert.glsl` so they share the same fragment shader. + +layout(set = 0, binding = 0, std140) uniform Globals { + mat4 transform; + uint input_colorspace; + uint _pad0; + uint _pad1; + uint _pad2; +} globals; + +layout(location = 0) in vec3 in_pos; +layout(location = 1) in vec4 in_color; +layout(location = 2) in vec2 in_uv; +layout(location = 3) in ivec2 in_layers; +layout(location = 4) in vec4 in_corner_radii; +layout(location = 5) in vec2 in_rect_size; +layout(location = 6) in int in_underline_style; +layout(location = 7) in vec4 in_clip_rect; + +layout(location = 0) out vec4 out_color; +layout(location = 1) out vec2 out_uv; +layout(location = 2) flat out ivec2 out_layers; +layout(location = 3) out vec4 out_corner_radii; +layout(location = 4) out vec2 out_rect_size; +layout(location = 5) flat out int out_underline_style; +layout(location = 6) flat out vec4 out_clip_rect; + +void main() { + gl_Position = globals.transform * vec4(in_pos, 1.0); + out_color = in_color; + out_uv = in_uv; + out_layers = in_layers; + out_corner_radii = in_corner_radii; + out_rect_size = in_rect_size; + out_underline_style = in_underline_style; + out_clip_rect = in_clip_rect; +} diff --git a/sugarloaf/src/renderer/image.frag.glsl b/sugarloaf/src/renderer/image.frag.glsl new file mode 100644 index 00000000..d25caf17 --- /dev/null +++ b/sugarloaf/src/renderer/image.frag.glsl @@ -0,0 +1,52 @@ +#version 450 + +// Image fragment shader, ported from `image_fs_main` in +// `sugarloaf/src/renderer/image.metal`. The texture is +// `R8G8B8A8_SRGB`, so the HW sRGB-decodes bytes on sample — +// bilinear filtering happens in *linear* light, matching Metal's +// `RGBA8Unorm_sRGB` behaviour. We then `linear_to_srgb` back to +// gamma encoding for the framebuffer (plain `B8G8R8A8_UNORM`, +// gamma-space alpha blending — same as the rest of sugarloaf). + +layout(set = 0, binding = 0, std140) uniform Globals { + mat4 transform; + uint input_colorspace; + uint _pad0; + uint _pad1; + uint _pad2; +} globals; + +layout(set = 1, binding = 0) uniform sampler2D image_texture; + +layout(location = 0) in vec2 in_tex_coord; + +layout(location = 0) out vec4 out_color; + +vec3 linear_to_srgb(vec3 c) { + vec3 lo = c * 12.92; + vec3 hi = pow(c, vec3(1.0 / 2.4)) * 1.055 - 0.055; + return mix(lo, hi, greaterThan(c, vec3(0.0031308))); +} + +vec3 rec2020_to_p3(vec3 linear_r2020) { + return vec3( + dot(linear_r2020, vec3( 1.34357825, -0.28217967, -0.06139858)), + dot(linear_r2020, vec3(-0.06529745, 1.08782226, -0.02252481)), + dot(linear_r2020, vec3( 0.00282179, -0.02598807, 1.02316628)) + ); +} + +void main() { + // R8G8B8A8_SRGB: HW decodes the byte → linear before handing it + // to the shader. Alpha is linear by convention regardless of + // format. After optional Rec.2020 → P3 gamut conversion, encode + // back to gamma-sRGB for the framebuffer. + vec4 rgba = texture(image_texture, in_tex_coord); + vec3 lin = rgba.rgb; + if (globals.input_colorspace == 2u) { + lin = rec2020_to_p3(lin); + } + vec3 enc = linear_to_srgb(lin); + // Premultiply alpha — pipeline blend factors are ONE / ONE_MINUS_SRC_ALPHA. + out_color = vec4(enc * rgba.a, rgba.a); +} diff --git a/sugarloaf/src/renderer/image.vert.glsl b/sugarloaf/src/renderer/image.vert.glsl new file mode 100644 index 00000000..71359f54 --- /dev/null +++ b/sugarloaf/src/renderer/image.vert.glsl @@ -0,0 +1,42 @@ +#version 450 + +// Image vertex shader, ported from `image_vs_main` in +// `sugarloaf/src/renderer/image.metal`. One instance per image +// placement (background image, kitty graphic, sixel image), 4 +// vertices per instance drawn as TRIANGLE_STRIP. +// +// Per-instance vertex layout matches `ImageInstance` in +// `sugarloaf/src/renderer/mod.rs` (32 bytes): +// loc 0 R32G32_SFLOAT dest_pos (offset 0) +// loc 1 R32G32_SFLOAT dest_size (offset 8) +// loc 2 R32G32B32A32_SFLOAT source_rect (offset 16) + +layout(set = 0, binding = 0, std140) uniform Globals { + mat4 transform; + uint input_colorspace; + uint _pad0; + uint _pad1; + uint _pad2; +} globals; + +layout(location = 0) in vec2 in_dest_pos; +layout(location = 1) in vec2 in_dest_size; +layout(location = 2) in vec4 in_source_rect; + +layout(location = 0) out vec2 out_tex_coord; + +void main() { + // Triangle strip: 4 vertices → quad + // 0 → 1 + // | /| + // 2 → 3 + vec2 corner; + corner.x = float(gl_VertexIndex == 1 || gl_VertexIndex == 3); + corner.y = float(gl_VertexIndex == 2 || gl_VertexIndex == 3); + + // `source_rect` is `[u0, v0, u1, v1]` (origin, end). + out_tex_coord = mix(in_source_rect.xy, in_source_rect.zw, corner); + + vec2 image_pos = in_dest_pos + in_dest_size * corner; + gl_Position = globals.transform * vec4(image_pos, 0.0, 1.0); +} diff --git a/sugarloaf/src/renderer/image_cache/atlas.rs b/sugarloaf/src/renderer/image_cache/atlas.rs index 29177023..09352a70 100644 --- a/sugarloaf/src/renderer/image_cache/atlas.rs +++ b/sugarloaf/src/renderer/image_cache/atlas.rs @@ -1,3 +1,6 @@ +// `grow_to` is reserved for atlas-resize on the rich-text path. +#![allow(dead_code)] + /// Improved shelf-based atlas allocator for better space utilization #[derive(Clone)] pub struct AtlasAllocator { diff --git a/sugarloaf/src/renderer/image_cache/cache.rs b/sugarloaf/src/renderer/image_cache/cache.rs index 8714c00d..486d3395 100644 --- a/sugarloaf/src/renderer/image_cache/cache.rs +++ b/sugarloaf/src/renderer/image_cache/cache.rs @@ -1,3 +1,9 @@ +// See `image_cache/mod.rs` — this whole module is reserved for the +// rich-text image-cache APIs that the grid renderer + UI text +// overlay are gradually superseding. Blanket-allow dead code rather +// than litter every reserved item with `#[allow]`. +#![allow(dead_code)] + use crate::context::{Context, ContextType}; use tracing::debug; @@ -74,6 +80,7 @@ struct ColorAtlasWithTexture { } enum ColorAtlasTexture { + #[cfg(feature = "wgpu")] Wgpu(wgpu::Texture, wgpu::TextureView), #[cfg(target_os = "macos")] Metal(metal::Texture), @@ -82,6 +89,7 @@ enum ColorAtlasTexture { } enum DeviceQueue { + #[cfg(feature = "wgpu")] Wgpu { device: std::sync::Arc, queue: std::sync::Arc, @@ -108,6 +116,9 @@ impl ImageCache { /// Creates a new image cache with mask atlas + initial color atlas pub fn new(context: &Context) -> Self { match &context.inner { + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), + #[cfg(feature = "wgpu")] ContextType::Wgpu(wgpu_context) => { let max_size = wgpu_context.max_texture_dimension_2d(); let max_texture_size = std::cmp::min(4096, max_size) as u16; @@ -232,6 +243,25 @@ impl ImageCache { device_queue: DeviceQueue::Cpu, } } + // Phase 1 stub: the Vulkan backend doesn't draw rich-text + // yet, so the image cache has nothing to upload. Mirror + // the CPU path — RAM-resident atlases, no GPU mirror. + // Phase 6 grows this into real Vulkan texture allocation. + #[cfg(target_os = "linux")] + ContextType::Vulkan(_) => { + let max_texture_size: u16 = 2048; + let color_atlases = vec![ColorAtlasWithTexture { + atlas: Atlas::new(AtlasKind::Color, max_texture_size), + texture: ColorAtlasTexture::Cpu, + }]; + Self { + entries: Vec::new(), + mask_atlas: Atlas::new(AtlasKind::Mask, max_texture_size), + color_atlases, + max_texture_size, + device_queue: DeviceQueue::Cpu, + } + } } } @@ -415,6 +445,7 @@ impl ImageCache { debug!("Creating color atlas {}", atlas_index); match &self.device_queue { + #[cfg(feature = "wgpu")] DeviceQueue::Wgpu { device, queue: _, .. } => { @@ -765,6 +796,9 @@ impl ImageCache { #[inline] pub fn process_atlases(&mut self, context: &mut Context) { match &context.inner { + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), + #[cfg(feature = "wgpu")] ContextType::Wgpu(wgpu_context) => { // Process mask atlas if self.mask_atlas.dirty { @@ -851,6 +885,17 @@ impl ImageCache { atlas_with_texture.atlas.dirty = false; } } + // Phase 1 stub: same as CPU — no GPU upload, atlas data + // lives in RAM. Phase 6 wires real Vulkan atlas uploads. + #[cfg(target_os = "linux")] + ContextType::Vulkan(_) => { + self.mask_atlas.fresh = false; + self.mask_atlas.dirty = false; + for atlas_with_texture in &mut self.color_atlases { + atlas_with_texture.atlas.fresh = false; + atlas_with_texture.atlas.dirty = false; + } + } #[cfg(target_os = "macos")] ContextType::Metal(_metal_context) => { // Process mask atlas @@ -911,6 +956,7 @@ impl ImageCache { } /// Get all texture views for WebGPU rendering (for texture array) + #[cfg(feature = "wgpu")] pub fn get_texture_views(&self) -> Vec<&wgpu::TextureView> { self.color_atlases .iter() @@ -940,6 +986,7 @@ impl ImageCache { } /// Get the mask texture view for WebGPU rendering + #[cfg(feature = "wgpu")] pub fn get_mask_texture_view(&self) -> Option<&wgpu::TextureView> { match &self.device_queue { DeviceQueue::Wgpu { diff --git a/sugarloaf/src/renderer/image_cache/mod.rs b/sugarloaf/src/renderer/image_cache/mod.rs index 1701fad7..ba98ae15 100644 --- a/sugarloaf/src/renderer/image_cache/mod.rs +++ b/sugarloaf/src/renderer/image_cache/mod.rs @@ -1,3 +1,10 @@ +// Lots of dead code in this module — held in reserve for the +// rich-text image-cache APIs that the grid renderer + UI text +// overlay are gradually superseding. Suppress the warnings rather +// than delete; the code is documented and will be revisited when +// the rich-text path is removed entirely. +#![allow(dead_code)] + pub(crate) mod atlas; mod cache; diff --git a/sugarloaf/src/renderer/mod.rs b/sugarloaf/src/renderer/mod.rs index 9426dd98..ecd58a45 100644 --- a/sugarloaf/src/renderer/mod.rs +++ b/sugarloaf/src/renderer/mod.rs @@ -2,8 +2,11 @@ mod batch; mod compositor; pub mod cpu; pub(crate) mod image_cache; +#[cfg(target_os = "linux")] +pub mod vulkan; use crate::components::core::orthographic_projection; +#[cfg(feature = "wgpu")] use crate::context::webgpu::WgpuContext; use crate::context::{Context, ContextType}; use crate::font::FontLibrary; @@ -19,6 +22,7 @@ use std::{borrow::Cow, mem}; #[cfg(target_os = "macos")] use parking_lot::Mutex; +#[cfg(feature = "wgpu")] use wgpu::util::DeviceExt; #[cfg(target_os = "macos")] @@ -26,6 +30,7 @@ use crate::context::metal::MetalContext; #[cfg(target_os = "macos")] use metal::*; +#[cfg(feature = "wgpu")] pub const BLEND: Option = Some(wgpu::BlendState { color: wgpu::BlendComponent { src_factor: wgpu::BlendFactor::SrcAlpha, @@ -47,13 +52,20 @@ pub const BLEND: Option = Some(wgpu::BlendState { // on every render call. #[allow(clippy::large_enum_variant)] pub enum RendererType { + #[cfg(feature = "wgpu")] Wgpu(WgpuRenderer), #[cfg(target_os = "macos")] Metal(MetalRenderer), + /// Native Vulkan backend (Linux). Mirrors the `Metal` variant in + /// scope: no librashader filters. Phase 1 = clear-and-present; + /// real pipelines land in later phases. + #[cfg(target_os = "linux")] + Vulkan(vulkan::VulkanRenderer), /// CPU backend: no GPU brush; rasterization happens in `cpu::CpuPipeline` at present time. Cpu, } +#[cfg(feature = "wgpu")] pub struct WgpuRenderer { vertex_buffer: wgpu::Buffer, instance_buffer: wgpu::Buffer, @@ -709,13 +721,26 @@ struct CachedGraphic { /// Backend-agnostic per-image GPU texture. /// Dropped when removed from the map → GPU memory freed immediately. +/// +/// The Vulkan variant carries an inline `VulkanImageTexture` +/// (image + view + memory + per-image descriptor pool/set) which is +/// larger than the other variants. Boxing it would buy nothing — +/// each entry is constructed once, lives until eviction, and is +/// only stored in the per-image map. +#[allow(clippy::large_enum_variant)] enum ImageTexture { + #[cfg(feature = "wgpu")] Wgpu { _texture: wgpu::Texture, // kept alive so `view` stays valid view: wgpu::TextureView, }, #[cfg(target_os = "macos")] Metal(metal::Texture), + /// Native Vulkan upload — owns image + view + descriptor set + /// (the descriptor set's binding 0 is wired to the image view + + /// shared sampler at upload time, so draw paths just bind it). + #[cfg(target_os = "linux")] + Vulkan(vulkan::VulkanImageTexture), } /// Per-image texture entry stored in the renderer. @@ -726,15 +751,20 @@ struct ImageTextureEntry { /// Per-instance data for image rendering (one instance = one image placement). /// The vertex shader generates 4 quad corners from vertex_id. +/// +/// `pub` because it appears in the signature of +/// `vulkan::VulkanRenderer::render_image_overlays` (also `pub` so +/// the `Renderer` dispatcher can call it). Not part of the crate's +/// public API in spirit — just in visibility. #[repr(C)] #[derive(Debug, Clone, Copy, bytemuck::Zeroable, bytemuck::Pod)] -struct ImageInstance { +pub struct ImageInstance { /// Screen position of the image top-left (physical pixels). - dest_pos: [f32; 2], + pub dest_pos: [f32; 2], /// Size of the image on screen (physical pixels). - dest_size: [f32; 2], + pub dest_size: [f32; 2], /// Source rectangle in the texture: xy = origin, zw = size (normalized 0..1). - source_rect: [f32; 4], + pub source_rect: [f32; 4], } /// Which layer to render the image in (relative to text). @@ -793,6 +823,26 @@ fn upload_background_image_texture( } let gpu = match &context.inner { crate::context::ContextType::Cpu(_) => return None, + // Vulkan path: the renderer owns the descriptor-set layout + // and shared sampler; we read them off the live brush_type + // here. Only the renderer is on the Sugarloaf struct, not + // the context, so the call site below threads them in. + // Actually: this function is a free fn taking only the + // context — we need to defer the upload until we have the + // renderer too. We do that by panicking here and pushing + // the real upload into a renderer method (see + // `Renderer::upload_background_image_vulkan`). When the + // dispatcher (`prepare`) sees a Vulkan ctx + dirty pixels, + // it calls the renderer method directly instead of this + // free function. + #[cfg(target_os = "linux")] + crate::context::ContextType::Vulkan(_) => { + unreachable!( + "Vulkan path uses Renderer::upload_background_image_vulkan, \ + not upload_background_image_texture" + ) + } + #[cfg(feature = "wgpu")] crate::context::ContextType::Wgpu(ctx) => { let texture = ctx.device.create_texture(&wgpu::TextureDescriptor { label: Some("sugarloaf::background image"), @@ -871,6 +921,8 @@ fn upload_background_image_texture( ); ImageTexture::Metal(mtl_tex) } + #[cfg(not(any(feature = "wgpu", target_os = "macos")))] + _ => return None, }; Some(ImageTextureEntry { gpu, @@ -886,6 +938,7 @@ impl Renderer { #[cfg(not(target_os = "macos"))] let _ = colorspace; let brush_type = match &context.inner { + #[cfg(feature = "wgpu")] ContextType::Wgpu(wgpu_context) => { RendererType::Wgpu(WgpuRenderer::new(wgpu_context)) } @@ -893,7 +946,13 @@ impl Renderer { ContextType::Metal(metal_context) => { RendererType::Metal(MetalRenderer::new(metal_context, colorspace)) } + #[cfg(target_os = "linux")] + ContextType::Vulkan(vulkan_context) => RendererType::Vulkan( + vulkan::VulkanRenderer::new(vulkan_context, colorspace), + ), ContextType::Cpu(_) => RendererType::Cpu, + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), }; Self { @@ -1126,8 +1185,24 @@ impl Renderer { // begins. The texture stays cached until a new image arrives or // `set_background_image_pixels(None)` is called. if let Some(pixels) = self.background_image_dirty.take() { - self.background_image_texture = - upload_background_image_texture(context, &pixels); + // Vulkan needs the renderer's descriptor-set layout + + // sampler to wire the per-image descriptor set, so it + // takes a different path that knows about both. + #[cfg(target_os = "linux")] + let used_vulkan = + if matches!(&context.inner, crate::context::ContextType::Vulkan(_)) { + self.upload_background_image_vulkan(context, &pixels); + true + } else { + false + }; + #[cfg(not(target_os = "linux"))] + let used_vulkan = false; + + if !used_vulkan { + self.background_image_texture = + upload_background_image_texture(context, &pixels); + } } self.instances.clear(); @@ -1220,10 +1295,9 @@ impl Renderer { } } + /// Render image overlays using per-image GPU textures. #[inline] #[allow(clippy::too_many_arguments)] - - /// Render image overlays using per-image GPU textures. fn render_graphic_overlays( &mut self, context: &mut crate::context::Context, @@ -1268,8 +1342,43 @@ impl Renderer { if matches!(&context.inner, crate::context::ContextType::Cpu(_)) { continue; } + // Vulkan: synchronous one-shot upload via the renderer's + // descriptor-set layout + sampler. The submit-and-wait + // is fine here — kitty placements come in bursts (a + // single image transmit, then many placements), and the + // upload is the cost we'd pay regardless. Move to a + // deferred per-frame pattern later if profiling shows + // image-heavy workloads stall. + #[cfg(target_os = "linux")] + if let crate::context::ContextType::Vulkan(vk_ctx) = &context.inner { + let RendererType::Vulkan(brush) = &self.brush_type else { + continue; + }; + let texture = vulkan::VulkanImageTexture::upload_rgba( + vk_ctx, + pixels, + width, + height, + brush.image_texture_descriptor_set_layout, + brush.image_sampler, + ); + self.image_textures.insert( + overlay.image_id, + ImageTextureEntry { + gpu: ImageTexture::Vulkan(texture), + transmit_time: entry.transmit_time, + }, + ); + continue; + } let gpu = match &context.inner { crate::context::ContextType::Cpu(_) => unreachable!(), + #[cfg(target_os = "linux")] + crate::context::ContextType::Vulkan(_) => unreachable!(), + #[cfg(not(feature = "wgpu"))] + #[allow(unreachable_patterns)] + _ => continue, + #[cfg(feature = "wgpu")] crate::context::ContextType::Wgpu(ctx) => { let texture = ctx.device.create_texture(&wgpu::TextureDescriptor { label: Some("kitty image"), @@ -1803,6 +1912,7 @@ impl Renderer { } #[inline] + #[cfg(feature = "wgpu")] pub fn render<'pass>( &'pass mut self, ctx: &mut WgpuContext, @@ -1821,7 +1931,6 @@ impl Renderer { .. } = self; - #[cfg_attr(not(target_os = "macos"), expect(irrefutable_let_patterns))] if let RendererType::Wgpu(brush) = brush_type { let color_views = images.get_texture_views(); let mask_texture_view = images.get_mask_texture_view(); @@ -2310,6 +2419,115 @@ impl Renderer { } } + /// Synchronously upload a background image into a Vulkan texture + /// and descriptor set. Called from the prepare path when the + /// user calls `Sugarloaf::set_background_image`. The + /// submit-and-wait is acceptable for one-shot uploads + /// (config-load time); kitty per-frame images take a different + /// deferred path. + #[cfg(target_os = "linux")] + fn upload_background_image_vulkan( + &mut self, + context: &crate::context::Context, + pixels: &BackgroundImagePixels, + ) { + let crate::context::ContextType::Vulkan(ctx) = &context.inner else { + return; + }; + let RendererType::Vulkan(brush) = &self.brush_type else { + return; + }; + let texture = vulkan::VulkanImageTexture::upload_rgba( + ctx, + &pixels.pixels, + pixels.width, + pixels.height, + brush.image_texture_descriptor_set_layout, + brush.image_sampler, + ); + self.background_image_texture = Some(ImageTextureEntry { + gpu: ImageTexture::Vulkan(texture), + transmit_time: std::time::Instant::now(), + }); + } + + /// Record sugarloaf's own draws inside the active dynamic-rendering + /// pass that `Sugarloaf::render_vulkan` opens. Order: + /// 1. Background image (full-screen quad). + /// 2. BelowText image overlays (kitty / sixel placements with + /// `dest_pos.z < 0`). + /// 3. Rich-text quad pass — `quad()` / `rect()` calls + cell + /// underline decorations (dashed/dotted/curly handled in + /// `quad.frag.glsl`). + /// 4. Non-quad geometry — `polygon()` / `line()` / `triangle()` + /// / `arc()` calls (cursor underline shape, hint highlights). + /// 5. AboveText image overlays. + /// 6. Optional bootstrap rect (`RIO_VULKAN_BOOTSTRAP=1`). + /// + /// Glyph atlas sampling through this pipeline isn't ported — + /// grid text + UI text overlay each own dedicated atlas + /// pipelines, so the rich-text path doesn't need it. + #[cfg(target_os = "linux")] + pub fn render_vulkan( + &mut self, + cmd: ash::vk::CommandBuffer, + frame: &crate::context::vulkan::VulkanFrame, + ) { + let viewport = [frame.extent.width as f32, frame.extent.height as f32]; + let slot = frame.slot; + + // Resolve image draws into (descriptor_set, instance) pairs + // before the &mut brush borrow takes hold; the per-image + // texture lookup needs an immutable borrow on + // `self.image_textures` which would conflict with the + // brush's `&mut self`. + let below: Vec<(ash::vk::DescriptorSet, ImageInstance)> = self + .image_draws + .iter() + .filter(|d| d.layer == ImageLayer::BelowText) + .filter_map(|d| { + let entry = self.image_textures.get(&d.image_id)?; + if let ImageTexture::Vulkan(tex) = &entry.gpu { + Some((tex.descriptor_set, d.instance)) + } else { + None + } + }) + .collect(); + let above: Vec<(ash::vk::DescriptorSet, ImageInstance)> = self + .image_draws + .iter() + .filter(|d| d.layer == ImageLayer::AboveText) + .filter_map(|d| { + let entry = self.image_textures.get(&d.image_id)?; + if let ImageTexture::Vulkan(tex) = &entry.gpu { + Some((tex.descriptor_set, d.instance)) + } else { + None + } + }) + .collect(); + + if let RendererType::Vulkan(brush) = &mut self.brush_type { + if let Some(bg) = &self.background_image_texture { + if let ImageTexture::Vulkan(tex) = &bg.gpu { + brush.render_background_image( + cmd, + slot, + viewport, + tex.descriptor_set, + ); + } + } + + brush.render_image_overlays(cmd, slot, viewport, &below); + brush.render_quads(cmd, slot, viewport, &self.instances); + brush.render_geometry(cmd, slot, viewport, &self.vertices); + brush.render_image_overlays(cmd, slot, viewport, &above); + brush.draw_bootstrap(cmd); + } + } + /// Vertices accumulated for the current frame (CPU rasterizer reads these). pub(crate) fn vertices(&self) -> &[Vertex] { &self.vertices @@ -2322,6 +2540,7 @@ impl Renderer { pub fn resize(&mut self, context: &mut Context) { let transform = match &context.inner { + #[cfg(feature = "wgpu")] ContextType::Wgpu(wgpu_ctx) => { orthographic_projection(wgpu_ctx.size.width, wgpu_ctx.size.height) } @@ -2329,12 +2548,19 @@ impl Renderer { ContextType::Metal(metal_ctx) => { orthographic_projection(metal_ctx.size.width, metal_ctx.size.height) } + #[cfg(target_os = "linux")] + ContextType::Vulkan(vulkan_ctx) => { + orthographic_projection(vulkan_ctx.size.width, vulkan_ctx.size.height) + } ContextType::Cpu(cpu_ctx) => { orthographic_projection(cpu_ctx.size.width, cpu_ctx.size.height) } + #[cfg(not(feature = "wgpu"))] + ContextType::_Phantom(_) => unreachable!(), }; match &mut self.brush_type { + #[cfg(feature = "wgpu")] RendererType::Wgpu(wgpu_brush) => { if transform != wgpu_brush.current_transform { let queue = match &context.inner { @@ -2359,11 +2585,19 @@ impl Renderer { // on the next frame's `orthographic_projection` call. let _ = transform; } + #[cfg(target_os = "linux")] + RendererType::Vulkan(_vulkan_brush) => { + // No-op: viewport + scissor are dynamic state set per + // frame in `VulkanRenderer::render`. The swapchain + // itself is rebuilt by `VulkanContext::resize`. + let _ = transform; + } RendererType::Cpu => {} } } } +#[cfg(feature = "wgpu")] impl WgpuRenderer { pub fn new(context: &WgpuContext) -> Self { let supported_vertex_buffer = 500; diff --git a/sugarloaf/src/renderer/quad.frag.glsl b/sugarloaf/src/renderer/quad.frag.glsl new file mode 100644 index 00000000..38d76b72 --- /dev/null +++ b/sugarloaf/src/renderer/quad.frag.glsl @@ -0,0 +1,232 @@ +#version 450 + +// Quad fragment shader — port of `fs_main` from +// `sugarloaf/src/renderer/renderer.metal`. Implements: +// * scissor-style clipping via `clip_rect` +// * underline pattern rasterization (regular / dashed / dotted / +// curly) — used by rich-text decorations and terminal cell +// underlines that emit through the rich-text quad path. +// * SDF rounded-corner mask (with edge antialiasing) +// * sRGB → output colorspace transform on the fill color +// +// NOT ported (yet): +// * glyph atlas sampling (color_layer / mask_layer paths) — the +// grid text pass and UI text overlay each own a dedicated +// atlas-sampling pipeline; the rich-text path doesn't render +// terminal glyphs. +// +// Fragments with `color_layer > 0` or `mask_layer > 0` (rich-text +// glyph instances that slip through here) fall back to the solid- +// fill path — graceful degradation rather than a visual crash. + +layout(set = 0, binding = 0, std140) uniform Globals { + mat4 transform; + uint input_colorspace; + uint _pad0; + uint _pad1; + uint _pad2; +} globals; + +layout(location = 0) in vec4 in_color; +layout(location = 1) in vec2 in_uv; +layout(location = 2) flat in ivec2 in_layers; +layout(location = 3) in vec4 in_corner_radii; +layout(location = 4) in vec2 in_rect_size; +layout(location = 5) flat in int in_underline_style; +layout(location = 6) flat in vec4 in_clip_rect; + +layout(location = 0) out vec4 out_color; + +// ----- Colorspace helpers (mirror grid_bg.frag.glsl) ----- +vec3 srgb_to_linear(vec3 c) { + vec3 lo = c / 12.92; + vec3 hi = pow((c + 0.055) / 1.055, vec3(2.4)); + return mix(lo, hi, greaterThan(c, vec3(0.04045))); +} +vec3 linear_to_srgb(vec3 c) { + vec3 lo = c * 12.92; + vec3 hi = pow(c, vec3(1.0 / 2.4)) * 1.055 - 0.055; + return mix(lo, hi, greaterThan(c, vec3(0.0031308))); +} +vec3 srgb_to_p3(vec3 linear_srgb) { + return vec3( + dot(linear_srgb, vec3(0.82246197, 0.17753803, 0.0)), + dot(linear_srgb, vec3(0.03319420, 0.96680580, 0.0)), + dot(linear_srgb, vec3(0.01708263, 0.07239744, 0.91051993)) + ); +} +vec3 rec2020_to_p3(vec3 linear_r2020) { + return vec3( + dot(linear_r2020, vec3( 1.34357825, -0.28217967, -0.06139858)), + dot(linear_r2020, vec3(-0.06529745, 1.08782226, -0.02252481)), + dot(linear_r2020, vec3( 0.00282179, -0.02598807, 1.02316628)) + ); +} +vec3 prepare_output_rgb(vec3 srgb, uint cs) { + vec3 lin = srgb_to_linear(srgb); + if (cs == 0u) { + lin = srgb_to_p3(lin); + } else if (cs == 2u) { + lin = rec2020_to_p3(lin); + } + return linear_to_srgb(lin); +} + +// ----- SDF rounded-rect helpers (mirror renderer.metal) ----- +float pick_corner_radius(vec2 center_to_point, vec4 corner_radii) { + if (center_to_point.x < 0.0) { + if (center_to_point.y < 0.0) { + return corner_radii.x; // tl + } else { + return corner_radii.w; // bl + } + } else { + if (center_to_point.y < 0.0) { + return corner_radii.y; // tr + } else { + return corner_radii.z; // br + } + } +} + +float quad_sdf(vec2 corner_center_to_point, float corner_radius) { + if (corner_radius == 0.0) { + return max(corner_center_to_point.x, corner_center_to_point.y); + } + float signed_distance_to_inset_quad = + length(max(vec2(0.0), corner_center_to_point)) + + min(0.0, max(corner_center_to_point.x, corner_center_to_point.y)); + return signed_distance_to_inset_quad - corner_radius; +} + +const float PI_F = 3.1415926; + +// Modulus that has the same sign as `a` (mirrors `fmod_pos` in +// renderer.metal — GLSL's `mod()` rounds toward negative infinity +// which is *not* what we want here). +float fmod_pos(float a, float b) { + return a - b * trunc(a / b); +} + +// Underline pattern alpha. Mirrors `underline_alpha` in +// renderer.metal. `x_pos`/`y_pos` are in pixels relative to the +// underline rect's top-left, `rect_height` is the rect's height in +// pixels, `thickness` is the line thickness encoded into +// `corner_radii.x` by the CPU emit code. `style` matches the +// `underline_style` enum from `batch.rs`: +// 1 = regular (solid), 2 = dashed, 3 = dotted, 4 = curly +float underline_alpha(float x_pos, float y_pos, float rect_height, + float thickness, int style) { + if (style == 1) { + return 1.0; + } + if (style == 2) { + // Dashed: 6px dash, 2px gap. + const float antialias = 0.5; + const float dash_width = 6.0; + const float gap_width = 2.0; + float period = dash_width + gap_width; + float pos_in_period = fmod_pos(x_pos, period); + float start_aa = clamp(pos_in_period / antialias, 0.0, 1.0); + float end_aa = clamp((dash_width - pos_in_period) / antialias, 0.0, 1.0); + return min(start_aa, end_aa); + } + if (style == 3) { + // Dotted: 2px dot, 2px gap. + const float antialias = 0.5; + const float dot_width = 2.0; + const float gap_width = 2.0; + float period = dot_width + gap_width; + float pos_in_period = fmod_pos(x_pos, period); + float start_aa = clamp(pos_in_period / antialias, 0.0, 1.0); + float end_aa = clamp((dot_width - pos_in_period) / antialias, 0.0, 1.0); + return min(start_aa, end_aa); + } + if (style == 4) { + // Curly (sine wave) via SDF. Same constants as renderer.metal. + const float WAVE_FREQUENCY = 2.0; + const float WAVE_HEIGHT_RATIO = 0.8; + + float half_thickness = thickness * 0.5; + vec2 st = vec2(x_pos / rect_height, y_pos / rect_height - 0.5); + float frequency = PI_F * WAVE_FREQUENCY * thickness / rect_height; + float amplitude = (thickness * WAVE_HEIGHT_RATIO) / rect_height; + + float sine = sin(st.x * frequency) * amplitude; + float dSine = cos(st.x * frequency) * amplitude * frequency; + float dist = (st.y - sine) / sqrt(1.0 + dSine * dSine); + float distance_in_pixels = dist * rect_height; + float distance_from_top_border = distance_in_pixels - half_thickness; + float distance_from_bottom_border = distance_in_pixels + half_thickness; + return clamp(0.5 - max(-distance_from_bottom_border, distance_from_top_border), + 0.0, 1.0); + } + return 1.0; +} + +void main() { + // Scissor-style clip rect. + if (in_clip_rect.z > 0.0) { + float px = gl_FragCoord.x; + float py = gl_FragCoord.y; + if (px < in_clip_rect.x + || px >= in_clip_rect.x + in_clip_rect.z + || py < in_clip_rect.y + || py >= in_clip_rect.y + in_clip_rect.w) + { + discard; + } + } + + vec4 color = in_color; + + // Underline patterns are emitted as their own quad instances + // (see `Compositor::push_underline` in `renderer/batch.rs`): + // color = underline color + // rect_size = (width, line_height) in pixels + // corner_radii.x = thickness + // underline_style ∈ {1, 2, 3, 4} + // The fragment computes the per-pattern alpha and returns the + // colorspace-transformed color * alpha. No SDF / corner masking + // applied to underline rects — they're flat strips. + if (in_underline_style > 0) { + float width = in_rect_size.x; + float rect_height = in_rect_size.y; + float x_pos = in_uv.x * width; + float y_pos = in_uv.y * rect_height; + float thickness = in_corner_radii.x; + float a = underline_alpha(x_pos, y_pos, rect_height, thickness, + in_underline_style); + out_color = vec4( + prepare_output_rgb(color.rgb, globals.input_colorspace), + color.a * a + ); + return; + } + + bool has_corners = in_corner_radii.x != 0.0 + || in_corner_radii.y != 0.0 + || in_corner_radii.z != 0.0 + || in_corner_radii.w != 0.0; + + // Fast path: sharp corners, no SDF. + if (!has_corners) { + out_color = vec4(prepare_output_rgb(color.rgb, globals.input_colorspace), color.a); + return; + } + + vec2 size = in_rect_size; + vec2 half_size = size / 2.0; + vec2 center_to_point = (in_uv - 0.5) * size; + float corner_radius = pick_corner_radius(center_to_point, in_corner_radii); + vec2 corner_to_point = abs(center_to_point) - half_size; + vec2 corner_center_to_point = corner_to_point + corner_radius; + float outer_sdf = quad_sdf(corner_center_to_point, corner_radius); + + const float antialias_threshold = 0.5; + if (outer_sdf >= antialias_threshold) { + discard; + } + float edge = clamp(antialias_threshold - outer_sdf, 0.0, 1.0); + out_color = vec4(prepare_output_rgb(color.rgb, globals.input_colorspace), color.a * edge); +} diff --git a/sugarloaf/src/renderer/quad.vert.glsl b/sugarloaf/src/renderer/quad.vert.glsl new file mode 100644 index 00000000..46bc0ab0 --- /dev/null +++ b/sugarloaf/src/renderer/quad.vert.glsl @@ -0,0 +1,71 @@ +#version 450 + +// Per-instance quad vertex shader, ported from `vs_instanced` in +// `sugarloaf/src/renderer/renderer.metal`. One `QuadInstance` per +// instance, vertex_id picks the corner. 4 vertices per instance, +// drawn as `TRIANGLE_STRIP`. +// +// Per-instance vertex layout matches `QuadInstance` in +// `sugarloaf/src/renderer/batch.rs` (96 bytes): +// loc 0 R32G32B32_SFLOAT pos[3] (offset 0) +// loc 1 R32G32B32A32_SFLOAT color[4] (offset 12) +// loc 2 R32G32B32A32_SFLOAT uv_rect[4] (offset 28) +// loc 3 R32G32_SINT layers[2] (offset 44) +// loc 4 R32G32_SFLOAT size[2] (offset 52) +// loc 5 R32G32B32A32_SFLOAT corner_radii[4] (offset 60) +// loc 6 R32_SINT underline_style (offset 76) +// loc 7 R32G32B32A32_SFLOAT clip_rect[4] (offset 80) +// +// This Vulkan port renders only the SDF rounded-rect fill path; the +// glyph atlas sampling and underline pattern paths from +// renderer.metal are not yet ported (grid text + UI text overlay +// have their own pipelines for those). + +layout(set = 0, binding = 0, std140) uniform Globals { + mat4 transform; + uint input_colorspace; + uint _pad0; + uint _pad1; + uint _pad2; +} globals; + +layout(location = 0) in vec3 in_pos; +layout(location = 1) in vec4 in_color; +layout(location = 2) in vec4 in_uv_rect; +layout(location = 3) in ivec2 in_layers; +layout(location = 4) in vec2 in_size; +layout(location = 5) in vec4 in_corner_radii; +layout(location = 6) in int in_underline_style; +layout(location = 7) in vec4 in_clip_rect; + +layout(location = 0) out vec4 out_color; +layout(location = 1) out vec2 out_uv; +layout(location = 2) flat out ivec2 out_layers; +layout(location = 3) out vec4 out_corner_radii; +layout(location = 4) out vec2 out_rect_size; +layout(location = 5) flat out int out_underline_style; +layout(location = 6) flat out vec4 out_clip_rect; + +// Unit quad corners for triangle strip: TL, BL, TR, BR. +const vec2 UNIT_QUAD[4] = vec2[]( + vec2(0.0, 0.0), + vec2(0.0, 1.0), + vec2(1.0, 0.0), + vec2(1.0, 1.0) +); + +void main() { + vec2 unit = UNIT_QUAD[gl_VertexIndex]; + vec2 pos = in_pos.xy + unit * in_size; + vec2 uv = mix(in_uv_rect.xy, in_uv_rect.zw, unit); + + gl_Position = globals.transform * vec4(pos, in_pos.z, 1.0); + + out_color = in_color; + out_uv = uv; + out_layers = in_layers; + out_corner_radii = in_corner_radii; + out_rect_size = in_size; + out_underline_style = in_underline_style; + out_clip_rect = in_clip_rect; +} diff --git a/sugarloaf/src/renderer/vulkan.rs b/sugarloaf/src/renderer/vulkan.rs new file mode 100644 index 00000000..6ae7fd7d --- /dev/null +++ b/sugarloaf/src/renderer/vulkan.rs @@ -0,0 +1,1672 @@ +// Copyright (c) 2023-present, Raphael Amorim. +// +// This source code is licensed under the MIT license found in the +// LICENSE file in the root directory of this source tree. + +//! Native Vulkan renderer (Phase 1: clear + bootstrap pipeline). +//! +//! Mirrors `MetalRenderer` in shape: holds compiled pipelines + per-frame +//! resources, exposes a `render` method that records draw commands into a +//! caller-supplied command buffer. Phase 1 only constructs the bootstrap +//! pipeline (centered debug rect) and clears the swapchain image. The +//! rich-text / image / grid pipelines come in later phases — see the +//! plan in `context/vulkan.rs`. +//! +//! The bootstrap rect is invisible by default; set +//! `RIO_VULKAN_BOOTSTRAP=1` in the environment to make it visible (a +//! centered magenta quad). The pipeline is always *constructed* either +//! way so any SPIR-V / pipeline-state validation errors surface +//! immediately at sugarloaf startup, not lazily on first frame with the +//! flag enabled. + +use ash::vk; + +use crate::context::vulkan::{ + allocate_host_visible_buffer_raw, VulkanBuffer, VulkanContext, VulkanFrame, + VulkanImage, FRAMES_IN_FLIGHT, +}; +use crate::renderer::batch::{QuadInstance, Vertex}; +use crate::renderer::ImageInstance; + +/// Compiled SPIR-V. Generated at build time from the matching +/// `.glsl` source by `sugarloaf/build.rs` (which shells out to +/// `glslc` or `glslangValidator`) and dropped into `OUT_DIR`. +/// Edit the `.glsl` file and rebuild — there's no manual recompile +/// step. Source files live in `sugarloaf/src/renderer/`. +const CLEAR_VERT_SPV: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/clear.vert.spv")); +const CLEAR_FRAG_SPV: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/clear.frag.spv")); +const QUAD_VERT_SPV: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/quad.vert.spv")); +const QUAD_FRAG_SPV: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/quad.frag.spv")); +const GEOMETRY_VERT_SPV: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/geometry.vert.spv")); +const IMAGE_VERT_SPV: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/image.vert.spv")); +const IMAGE_FRAG_SPV: &[u8] = include_bytes!(concat!(env!("OUT_DIR"), "/image.frag.spv")); + +/// `std140`-padded `Globals` uniform — `mat4 transform` (64 B) + +/// `uint input_colorspace` + 12 B padding to round up to 80 B (a +/// multiple of 16). Mirrors the `Globals` block in +/// `quad.{vert,frag}.glsl`. +#[repr(C)] +#[derive(Copy, Clone)] +struct QuadGlobals { + transform: [f32; 16], + input_colorspace: u32, + _pad: [u32; 3], +} + +pub struct VulkanRenderer { + /// Bootstrap pipeline — proves the SPIR-V → `vk::ShaderModule` → + /// `vk::Pipeline` chain works end-to-end. Bound only when + /// `bootstrap_visible` is true; otherwise the frame is just a + /// clear-and-present. + bootstrap_pipeline: vk::Pipeline, + bootstrap_layout: vk::PipelineLayout, + bootstrap_visible: bool, + /// Cloned from `VulkanContext::device()` so `Drop` can destroy our + /// pipelines even after the parent context has dropped its borrow. + /// `ash::Device` is `Clone` (just a wrapper around fn pointers + a + /// `vk::Device` handle); cloning does not create a new logical + /// device. + device: ash::Device, + /// Cached for buffer (re)allocation in `render_quads` (which only + /// has `&mut self`, no `&VulkanContext`). + instance: ash::Instance, + physical_device: vk::PhysicalDevice, + /// Set once at construction from the configured `Colorspace`. + /// Mirrors `MetalRenderer::input_colorspace`. Value is `0 = sRGB`, + /// `1 = DisplayP3`, `2 = Rec.2020`. + input_colorspace: u32, + + // ---------- quad pipeline (rich-text rect/rounded-rect path) ---------- + quad_pipeline: vk::Pipeline, + quad_pipeline_layout: vk::PipelineLayout, + quad_descriptor_pool: vk::DescriptorPool, + quad_descriptor_set_layout: vk::DescriptorSetLayout, + quad_descriptor_sets: [vk::DescriptorSet; FRAMES_IN_FLIGHT], + /// Per-slot `Globals` uniform buffer. + quad_uniform_buffers: [VulkanBuffer; FRAMES_IN_FLIGHT], + /// Per-slot per-instance `QuadInstance` vertex buffer ring. + /// Allocated lazily, grown on demand. + quad_instance_buffers: [Option; FRAMES_IN_FLIGHT], + quad_instance_capacity: [usize; FRAMES_IN_FLIGHT], + + // ---------- non-quad geometry pipeline ---------- + /// Renders `Vertex`-supplied geometry (`polygon()` / `line()` / + /// `triangle()` / `arc()` calls). Shares the quad pipeline's + /// descriptor set layout + uniform buffers + fragment shader; + /// only the vertex shader and topology differ. + geometry_pipeline: vk::Pipeline, + /// Per-slot per-vertex `Vertex` buffer ring. + geometry_vertex_buffers: [Option; FRAMES_IN_FLIGHT], + geometry_vertex_capacity: [usize; FRAMES_IN_FLIGHT], + + // ---------- image pipeline ---------- + /// Set 0 = `Globals` uniform (per slot, shared with quad pipeline + /// shape but separate buffers — image's `Globals` doesn't change + /// per draw and could be coalesced; for now just duplicate to + /// keep the descriptor sets simple). + /// Set 1 = single combined image+sampler. Owned per + /// `VulkanImageTexture` so each image carries its own descriptor. + image_pipeline: vk::Pipeline, + image_pipeline_layout: vk::PipelineLayout, + image_uniform_descriptor_set_layout: vk::DescriptorSetLayout, + /// Public so per-image `VulkanImageTexture` instances can allocate + /// their own descriptor sets at upload time. + pub image_texture_descriptor_set_layout: vk::DescriptorSetLayout, + image_uniform_descriptor_pool: vk::DescriptorPool, + image_uniform_descriptor_sets: [vk::DescriptorSet; FRAMES_IN_FLIGHT], + image_uniform_buffers: [VulkanBuffer; FRAMES_IN_FLIGHT], + /// Per-slot per-instance `ImageInstance` vertex buffer ring (for + /// kitty/sixel images). The bg image gets its own dedicated + /// 1-instance vertex buffer (`image_bg_vertex_buffers`). + image_instance_buffers: [Option; FRAMES_IN_FLIGHT], + image_instance_capacity: [usize; FRAMES_IN_FLIGHT], + /// Dedicated single-instance vertex buffer for the background + /// image (one per slot). Kept separate so the bg draw never + /// collides with kitty placement slots — same pattern as the + /// wgpu `background_image_vertex_buffer`. + image_bg_vertex_buffers: [VulkanBuffer; FRAMES_IN_FLIGHT], + /// Sampler shared by every image draw. Linear filtering for + /// background images (smooth scaling); kitty graphics also looks + /// fine with linear. + pub image_sampler: vk::Sampler, +} + +impl VulkanRenderer { + pub fn new(ctx: &VulkanContext, colorspace: crate::sugarloaf::Colorspace) -> Self { + let device = ctx.device().clone(); + let instance = ctx.instance().clone(); + let physical_device = ctx.physical_device(); + let color_format = ctx.swapchain_format(); + let input_colorspace = match colorspace { + crate::sugarloaf::Colorspace::Srgb => 0u32, + crate::sugarloaf::Colorspace::DisplayP3 => 1u32, + crate::sugarloaf::Colorspace::Rec2020 => 2u32, + }; + + let vert_module = create_shader_module(&device, CLEAR_VERT_SPV); + let frag_module = create_shader_module(&device, CLEAR_FRAG_SPV); + + // Push constant: vec4 color, fragment stage, offset 0 — matches + // `layout(push_constant) uniform PC { vec4 color; }` in + // `clear.frag.glsl`. + let push_constant_range = vk::PushConstantRange::default() + .stage_flags(vk::ShaderStageFlags::FRAGMENT) + .offset(0) + .size(std::mem::size_of::<[f32; 4]>() as u32); + + let push_constant_ranges = [push_constant_range]; + let layout_info = vk::PipelineLayoutCreateInfo::default() + .push_constant_ranges(&push_constant_ranges); + let bootstrap_layout = unsafe { + device + .create_pipeline_layout(&layout_info, None) + .expect("create_pipeline_layout(bootstrap)") + }; + + let pipeline_cache = ctx.pipeline_cache(); + let bootstrap_pipeline = build_bootstrap_pipeline( + &device, + pipeline_cache, + bootstrap_layout, + vert_module, + frag_module, + color_format, + ); + + // Shader modules are no longer needed once the pipeline is built. + // The compiled SPIR-V is baked into the pipeline state. + unsafe { + device.destroy_shader_module(vert_module, None); + device.destroy_shader_module(frag_module, None); + } + + let bootstrap_visible = std::env::var_os("RIO_VULKAN_BOOTSTRAP") + .map(|v| v != "0" && !v.is_empty()) + .unwrap_or(false); + + if bootstrap_visible { + tracing::info!("Vulkan bootstrap rect enabled (RIO_VULKAN_BOOTSTRAP set)"); + } + + // Quad pipeline construction. + let quad_uniform_buffers = std::array::from_fn(|_| { + ctx.allocate_host_visible_buffer( + std::mem::size_of::() as u64, + vk::BufferUsageFlags::UNIFORM_BUFFER, + ) + }); + let quad_descriptor_set_layout = create_quad_descriptor_set_layout(&device); + let quad_descriptor_pool = create_quad_descriptor_pool(&device); + let quad_descriptor_sets = allocate_quad_descriptor_sets( + &device, + quad_descriptor_pool, + quad_descriptor_set_layout, + ); + for slot in 0..FRAMES_IN_FLIGHT { + update_quad_descriptor_set( + &device, + quad_descriptor_sets[slot], + &quad_uniform_buffers[slot], + ); + } + let quad_pipeline_layout = + create_quad_pipeline_layout(&device, quad_descriptor_set_layout); + let quad_pipeline = build_quad_pipeline( + &device, + pipeline_cache, + quad_pipeline_layout, + color_format, + ); + // Geometry pipeline shares quad's descriptor set layout + + // uniform buffers + fragment shader; just a different + // vertex shader + input layout + topology. + let geometry_pipeline = build_geometry_pipeline( + &device, + pipeline_cache, + quad_pipeline_layout, + color_format, + ); + + // Image pipeline construction. + let image_uniform_descriptor_set_layout = + create_image_uniform_descriptor_set_layout(&device); + let image_texture_descriptor_set_layout = + create_image_texture_descriptor_set_layout(&device); + let image_uniform_descriptor_pool = create_image_uniform_descriptor_pool(&device); + let image_uniform_descriptor_sets = allocate_quad_descriptor_sets( + &device, + image_uniform_descriptor_pool, + image_uniform_descriptor_set_layout, + ); + let image_uniform_buffers = std::array::from_fn(|_| { + ctx.allocate_host_visible_buffer( + std::mem::size_of::() as u64, + vk::BufferUsageFlags::UNIFORM_BUFFER, + ) + }); + for slot in 0..FRAMES_IN_FLIGHT { + update_quad_descriptor_set( + &device, + image_uniform_descriptor_sets[slot], + &image_uniform_buffers[slot], + ); + } + let image_pipeline_layout = create_image_pipeline_layout( + &device, + image_uniform_descriptor_set_layout, + image_texture_descriptor_set_layout, + ); + let image_pipeline = build_image_pipeline( + &device, + pipeline_cache, + image_pipeline_layout, + color_format, + ); + + let image_bg_vertex_buffers = std::array::from_fn(|_| { + ctx.allocate_host_visible_buffer( + std::mem::size_of::() as u64, + vk::BufferUsageFlags::VERTEX_BUFFER, + ) + }); + let image_sampler = create_image_sampler(&device); + + Self { + bootstrap_pipeline, + bootstrap_layout, + bootstrap_visible, + device, + instance, + physical_device, + input_colorspace, + quad_pipeline, + quad_pipeline_layout, + quad_descriptor_pool, + quad_descriptor_set_layout, + quad_descriptor_sets, + quad_uniform_buffers, + quad_instance_buffers: std::array::from_fn(|_| None), + quad_instance_capacity: [0; FRAMES_IN_FLIGHT], + geometry_pipeline, + geometry_vertex_buffers: std::array::from_fn(|_| None), + geometry_vertex_capacity: [0; FRAMES_IN_FLIGHT], + image_pipeline, + image_pipeline_layout, + image_uniform_descriptor_set_layout, + image_texture_descriptor_set_layout, + image_uniform_descriptor_pool, + image_uniform_descriptor_sets, + image_uniform_buffers, + image_instance_buffers: std::array::from_fn(|_| None), + image_instance_capacity: [0; FRAMES_IN_FLIGHT], + image_bg_vertex_buffers, + image_sampler, + } + } + + /// Record the non-quad geometry pass into `cmd`. Caller has + /// already opened the dynamic-rendering pass and set + /// viewport/scissor. No-op when `vertices` is empty. + /// + /// Reuses `quad_pipeline_layout` + per-slot quad uniform set — + /// the Globals binding is the same shape; we already wrote it + /// in `render_quads` if there were quads this frame. If quads + /// were skipped, we have to upload here too. (Doing it + /// unconditionally keeps the call sites independent.) + pub fn render_geometry( + &mut self, + cmd: vk::CommandBuffer, + slot: usize, + viewport: [f32; 2], + vertices: &[Vertex], + ) { + if vertices.is_empty() { + return; + } + debug_assert!(slot < FRAMES_IN_FLIGHT); + + // Make sure the per-slot quad uniform buffer (shared with + // the geometry pipeline) holds the current frame's transform + // even when no quads were drawn. + let transform = + crate::components::core::orthographic_projection(viewport[0], viewport[1]); + let globals = QuadGlobals { + transform, + input_colorspace: self.input_colorspace, + _pad: [0; 3], + }; + unsafe { + let dst = self.quad_uniform_buffers[slot].as_mut_ptr() as *mut QuadGlobals; + std::ptr::write(dst, globals); + } + + // Grow per-slot vertex buffer if needed. + let vertex_count = vertices.len(); + let needed_bytes = std::mem::size_of_val(vertices); + if vertex_count > self.geometry_vertex_capacity[slot] { + let new_cap = vertex_count.next_power_of_two().max(256); + self.geometry_vertex_buffers[slot] = Some(allocate_host_visible_buffer_raw( + &self.device, + &self.instance, + self.physical_device, + (new_cap * std::mem::size_of::()) as u64, + vk::BufferUsageFlags::VERTEX_BUFFER, + )); + self.geometry_vertex_capacity[slot] = new_cap; + } + let vertex_buf = self.geometry_vertex_buffers[slot].as_ref().unwrap(); + unsafe { + std::ptr::copy_nonoverlapping( + vertices.as_ptr() as *const u8, + vertex_buf.as_mut_ptr(), + needed_bytes, + ); + } + + unsafe { + self.device.cmd_bind_pipeline( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.geometry_pipeline, + ); + self.device.cmd_bind_descriptor_sets( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.quad_pipeline_layout, + 0, + &[self.quad_descriptor_sets[slot]], + &[], + ); + self.device + .cmd_bind_vertex_buffers(cmd, 0, &[vertex_buf.handle()], &[0]); + // Caller-provided vertices are TRIANGLE_LIST — the emit + // path tessellates polygons / arcs / lines into + // triangles before pushing. + self.device.cmd_draw(cmd, vertex_count as u32, 1, 0, 0); + } + } + + /// Draw a batch of image overlays (kitty / sixel placements) for + /// one layer (BelowText or AboveText). Each `(descriptor_set, + /// instance)` pair is one image placement — caller has resolved + /// the per-image descriptor set ahead of time. Writes all + /// instances into the per-slot ring buffer in order, then issues + /// one `cmd_draw(4, 1, ...)` per placement, binding the buffer + /// at the matching byte offset. + /// + /// Pre-binds the uniform set + image pipeline once, then loops + /// per-image just to update the texture descriptor + vertex + /// buffer offset. + pub fn render_image_overlays( + &mut self, + cmd: vk::CommandBuffer, + slot: usize, + viewport: [f32; 2], + draws: &[(vk::DescriptorSet, ImageInstance)], + ) { + if draws.is_empty() { + return; + } + debug_assert!(slot < FRAMES_IN_FLIGHT); + + // Update the (shared) image uniform with the current + // viewport's transform. + let transform = + crate::components::core::orthographic_projection(viewport[0], viewport[1]); + let globals = QuadGlobals { + transform, + input_colorspace: self.input_colorspace, + _pad: [0; 3], + }; + unsafe { + let dst = self.image_uniform_buffers[slot].as_mut_ptr() as *mut QuadGlobals; + std::ptr::write(dst, globals); + } + + // Grow the per-slot kitty/sixel instance buffer if needed. + let count = draws.len(); + let stride = std::mem::size_of::(); + let needed_bytes = count * stride; + if count > self.image_instance_capacity[slot] { + let new_cap = count.next_power_of_two().max(16); + self.image_instance_buffers[slot] = Some(allocate_host_visible_buffer_raw( + &self.device, + &self.instance, + self.physical_device, + (new_cap * stride) as u64, + vk::BufferUsageFlags::VERTEX_BUFFER, + )); + self.image_instance_capacity[slot] = new_cap; + } + let buf = self.image_instance_buffers[slot].as_ref().unwrap(); + unsafe { + // Write instances contiguously, ordered the same as `draws`. + let dst = buf.as_mut_ptr() as *mut ImageInstance; + for (i, (_set, inst)) in draws.iter().enumerate() { + std::ptr::write(dst.add(i), *inst); + } + } + + unsafe { + self.device.cmd_bind_pipeline( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.image_pipeline, + ); + // Set 0 (uniform) is constant across draws — bind once. + self.device.cmd_bind_descriptor_sets( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.image_pipeline_layout, + 0, + &[self.image_uniform_descriptor_sets[slot]], + &[], + ); + for (i, (texture_set, _inst)) in draws.iter().enumerate() { + // Set 1 (texture) changes per-draw — rebind. + self.device.cmd_bind_descriptor_sets( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.image_pipeline_layout, + 1, + &[*texture_set], + &[], + ); + let byte_offset = (i * stride) as u64; + self.device.cmd_bind_vertex_buffers( + cmd, + 0, + &[buf.handle()], + &[byte_offset], + ); + let _ = needed_bytes; // shut up unused warning + self.device.cmd_draw(cmd, 4, 1, 0, 0); + } + } + } + + /// Draw the background image, if any. Caller passes the per-image + /// descriptor set (allocated at upload time via + /// `VulkanImageTexture::new`). Writes a single ImageInstance with + /// `dest_pos = [0, 0]` and `dest_size = viewport` so the image + /// covers the whole window. + pub fn render_background_image( + &mut self, + cmd: vk::CommandBuffer, + slot: usize, + viewport: [f32; 2], + image_texture_descriptor_set: vk::DescriptorSet, + ) { + debug_assert!(slot < FRAMES_IN_FLIGHT); + + // Update uniforms (transform). + let transform = + crate::components::core::orthographic_projection(viewport[0], viewport[1]); + let globals = QuadGlobals { + transform, + input_colorspace: self.input_colorspace, + _pad: [0; 3], + }; + unsafe { + let dst = self.image_uniform_buffers[slot].as_mut_ptr() as *mut QuadGlobals; + std::ptr::write(dst, globals); + } + + // Build a single full-screen ImageInstance. + let instance = ImageInstance { + dest_pos: [0.0, 0.0], + dest_size: viewport, + source_rect: [0.0, 0.0, 1.0, 1.0], + }; + unsafe { + let dst = + self.image_bg_vertex_buffers[slot].as_mut_ptr() as *mut ImageInstance; + std::ptr::write(dst, instance); + } + + unsafe { + self.device.cmd_bind_pipeline( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.image_pipeline, + ); + self.device.cmd_bind_descriptor_sets( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.image_pipeline_layout, + 0, + &[ + self.image_uniform_descriptor_sets[slot], + image_texture_descriptor_set, + ], + &[], + ); + self.device.cmd_bind_vertex_buffers( + cmd, + 0, + &[self.image_bg_vertex_buffers[slot].handle()], + &[0], + ); + self.device.cmd_draw(cmd, 4, 1, 0, 0); + } + } + + /// Record the rich-text quad pass into `cmd`. Caller has already + /// opened the dynamic-rendering pass and set viewport/scissor. + /// No-op when `instances` is empty. + pub fn render_quads( + &mut self, + cmd: vk::CommandBuffer, + slot: usize, + viewport: [f32; 2], + instances: &[QuadInstance], + ) { + if instances.is_empty() { + return; + } + debug_assert!(slot < FRAMES_IN_FLIGHT); + + // Update uniforms. + let transform = + crate::components::core::orthographic_projection(viewport[0], viewport[1]); + let globals = QuadGlobals { + transform, + input_colorspace: self.input_colorspace, + _pad: [0; 3], + }; + unsafe { + let dst = self.quad_uniform_buffers[slot].as_mut_ptr() as *mut QuadGlobals; + std::ptr::write(dst, globals); + } + + // Grow per-slot instance buffer if needed. + let instance_count = instances.len(); + let needed_bytes = std::mem::size_of_val(instances); + if instance_count > self.quad_instance_capacity[slot] { + let new_cap = instance_count.next_power_of_two().max(256); + self.quad_instance_buffers[slot] = Some(allocate_host_visible_buffer_raw( + &self.device, + &self.instance, + self.physical_device, + (new_cap * std::mem::size_of::()) as u64, + vk::BufferUsageFlags::VERTEX_BUFFER, + )); + self.quad_instance_capacity[slot] = new_cap; + } + let instance_buf = self.quad_instance_buffers[slot].as_ref().unwrap(); + unsafe { + std::ptr::copy_nonoverlapping( + instances.as_ptr() as *const u8, + instance_buf.as_mut_ptr(), + needed_bytes, + ); + } + + unsafe { + self.device.cmd_bind_pipeline( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.quad_pipeline, + ); + self.device.cmd_bind_descriptor_sets( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.quad_pipeline_layout, + 0, + &[self.quad_descriptor_sets[slot]], + &[], + ); + self.device + .cmd_bind_vertex_buffers(cmd, 0, &[instance_buf.handle()], &[0]); + // 4 vertices per instance (TRIANGLE_STRIP quad). + self.device.cmd_draw(cmd, 4, instance_count as u32, 0, 0); + } + } + + /// Whether the user opted into the magenta debug rect via + /// `RIO_VULKAN_BOOTSTRAP=1`. Read by `Sugarloaf::render_vulkan` + /// so the rect can be drawn between grid passes and the present + /// barrier — keeping all draws inside the single render pass that + /// the Sugarloaf-level orchestrator opens. + #[inline] + pub fn bootstrap_visible(&self) -> bool { + self.bootstrap_visible + } + + /// Record the bootstrap rect draw into `cmd`. Caller must already + /// have `cmd_begin_rendering` open + viewport/scissor set. No-op + /// when `bootstrap_visible == false`. + pub fn draw_bootstrap(&self, cmd: vk::CommandBuffer) { + if !self.bootstrap_visible { + return; + } + unsafe { + self.device.cmd_bind_pipeline( + cmd, + vk::PipelineBindPoint::GRAPHICS, + self.bootstrap_pipeline, + ); + let color: [f32; 4] = [1.0, 0.0, 1.0, 1.0]; + self.device.cmd_push_constants( + cmd, + self.bootstrap_layout, + vk::ShaderStageFlags::FRAGMENT, + 0, + bytemuck::bytes_of(&color), + ); + // Triangle strip, 4 vertices — `clear.vert.glsl` + // generates a centered rect in NDC. + self.device.cmd_draw(cmd, 4, 1, 0, 0); + } + } +} + +/// Free helper: emit the layout-transition barrier from `UNDEFINED` to +/// `COLOR_ATTACHMENT_OPTIMAL`. Run once at the top of a frame, before +/// `cmd_begin_rendering`. The discard of previous contents is +/// intentional — sugarloaf clears every frame, so we don't need to +/// preserve what the swapchain image held last present. +pub fn cmd_acquire_image_for_rendering( + device: &ash::Device, + cmd: vk::CommandBuffer, + image: vk::Image, +) { + unsafe { + let barrier = vk::ImageMemoryBarrier2::default() + .src_stage_mask(vk::PipelineStageFlags2::TOP_OF_PIPE) + .src_access_mask(vk::AccessFlags2::empty()) + .dst_stage_mask(vk::PipelineStageFlags2::COLOR_ATTACHMENT_OUTPUT) + .dst_access_mask(vk::AccessFlags2::COLOR_ATTACHMENT_WRITE) + .old_layout(vk::ImageLayout::UNDEFINED) + .new_layout(vk::ImageLayout::COLOR_ATTACHMENT_OPTIMAL) + .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .image(image) + .subresource_range(color_subresource_range()); + let barriers = [barrier]; + let dep = vk::DependencyInfo::default().image_memory_barriers(&barriers); + device.cmd_pipeline_barrier2(cmd, &dep); + } +} + +/// Free helper: emit the post-rendering barrier so +/// `vkQueuePresentKHR` finds the image in `PRESENT_SRC_KHR`. Run once +/// after `cmd_end_rendering`, before `present_frame`. +pub fn cmd_release_image_to_present( + device: &ash::Device, + cmd: vk::CommandBuffer, + image: vk::Image, +) { + unsafe { + let barrier = vk::ImageMemoryBarrier2::default() + .src_stage_mask(vk::PipelineStageFlags2::COLOR_ATTACHMENT_OUTPUT) + .src_access_mask(vk::AccessFlags2::COLOR_ATTACHMENT_WRITE) + .dst_stage_mask(vk::PipelineStageFlags2::BOTTOM_OF_PIPE) + .dst_access_mask(vk::AccessFlags2::empty()) + .old_layout(vk::ImageLayout::COLOR_ATTACHMENT_OPTIMAL) + .new_layout(vk::ImageLayout::PRESENT_SRC_KHR) + .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .image(image) + .subresource_range(color_subresource_range()); + let barriers = [barrier]; + let dep = vk::DependencyInfo::default().image_memory_barriers(&barriers); + device.cmd_pipeline_barrier2(cmd, &dep); + } +} + +/// Free helper: build a `RenderingInfo` clearing the swapchain image to +/// `clear_color`. Caller wraps draws between `cmd_begin_rendering` and +/// `cmd_end_rendering` using this. +pub fn build_rendering_info<'a>( + frame: &'a VulkanFrame, + color_attachments: &'a [vk::RenderingAttachmentInfo<'a>], +) -> vk::RenderingInfo<'a> { + let render_area = vk::Rect2D { + offset: vk::Offset2D { x: 0, y: 0 }, + extent: frame.extent, + }; + vk::RenderingInfo::default() + .render_area(render_area) + .layer_count(1) + .color_attachments(color_attachments) +} + +/// Build the single color attachment used by every sugarloaf Vulkan +/// frame: clear → store, no MSAA, no depth. +pub fn build_color_attachment( + frame: &VulkanFrame, + clear_color: [f32; 4], +) -> vk::RenderingAttachmentInfo<'_> { + let clear = vk::ClearValue { + color: vk::ClearColorValue { + float32: clear_color, + }, + }; + vk::RenderingAttachmentInfo::default() + .image_view(frame.image_view) + .image_layout(vk::ImageLayout::COLOR_ATTACHMENT_OPTIMAL) + .load_op(vk::AttachmentLoadOp::CLEAR) + .store_op(vk::AttachmentStoreOp::STORE) + .clear_value(clear) +} + +impl Drop for VulkanRenderer { + fn drop(&mut self) { + unsafe { + // Idle the device before tearing down pipelines. The parent + // `VulkanContext::Drop` also waits, but `Sugarloaf`'s field + // declaration order has us dropping first — and an outstanding + // submit using this pipeline would crash the driver if we + // destroyed it under the GPU's nose. + let _ = self.device.device_wait_idle(); + + self.device.destroy_pipeline(self.image_pipeline, None); + self.device + .destroy_pipeline_layout(self.image_pipeline_layout, None); + self.device.destroy_sampler(self.image_sampler, None); + self.device + .destroy_descriptor_pool(self.image_uniform_descriptor_pool, None); + self.device.destroy_descriptor_set_layout( + self.image_uniform_descriptor_set_layout, + None, + ); + self.device.destroy_descriptor_set_layout( + self.image_texture_descriptor_set_layout, + None, + ); + + self.device.destroy_pipeline(self.geometry_pipeline, None); + self.device.destroy_pipeline(self.quad_pipeline, None); + self.device + .destroy_pipeline_layout(self.quad_pipeline_layout, None); + self.device + .destroy_descriptor_pool(self.quad_descriptor_pool, None); + self.device + .destroy_descriptor_set_layout(self.quad_descriptor_set_layout, None); + self.device.destroy_pipeline(self.bootstrap_pipeline, None); + self.device + .destroy_pipeline_layout(self.bootstrap_layout, None); + // Buffers (uniform, instance) drop themselves via VulkanBuffer. + } + } +} + +// ----------------------------------------------------------------------- +// Internal helpers +// ----------------------------------------------------------------------- + +fn color_subresource_range() -> vk::ImageSubresourceRange { + vk::ImageSubresourceRange::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .base_mip_level(0) + .level_count(1) + .base_array_layer(0) + .layer_count(1) +} + +/// Wrap `bytes` (an `include_bytes!` slice) in a `vk::ShaderModule`. The +/// SPIR-V spec requires u32 alignment, but `include_bytes!` is +/// byte-aligned — `ash::util::read_spv` does the alignment-safe copy +/// for us. +fn create_shader_module(device: &ash::Device, bytes: &[u8]) -> vk::ShaderModule { + let code = ash::util::read_spv(&mut std::io::Cursor::new(bytes)) + .expect("read_spv (embedded shader is valid)"); + let info = vk::ShaderModuleCreateInfo::default().code(&code); + unsafe { + device + .create_shader_module(&info, None) + .expect("create_shader_module") + } +} + +fn build_bootstrap_pipeline( + device: &ash::Device, + pipeline_cache: vk::PipelineCache, + layout: vk::PipelineLayout, + vert: vk::ShaderModule, + frag: vk::ShaderModule, + color_format: vk::Format, +) -> vk::Pipeline { + let entry = c"main"; + let stages = [ + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::VERTEX) + .module(vert) + .name(entry), + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::FRAGMENT) + .module(frag) + .name(entry), + ]; + + // No vertex input — `clear.vert.glsl` uses `gl_VertexIndex` to + // derive positions, no buffers bound. + let vertex_input = vk::PipelineVertexInputStateCreateInfo::default(); + + let input_assembly = vk::PipelineInputAssemblyStateCreateInfo::default() + .topology(vk::PrimitiveTopology::TRIANGLE_STRIP) + .primitive_restart_enable(false); + + // Viewport + scissor are dynamic so resize doesn't need a pipeline + // rebuild. + let viewport_state = vk::PipelineViewportStateCreateInfo::default() + .viewport_count(1) + .scissor_count(1); + + let rasterization = vk::PipelineRasterizationStateCreateInfo::default() + .polygon_mode(vk::PolygonMode::FILL) + .cull_mode(vk::CullModeFlags::NONE) + .front_face(vk::FrontFace::COUNTER_CLOCKWISE) + .line_width(1.0); + + let multisample = vk::PipelineMultisampleStateCreateInfo::default() + .rasterization_samples(vk::SampleCountFlags::TYPE_1); + + // Premultiplied-over blend, matching the rest of the sugarloaf + // pipelines (Metal/wgpu compositor blend). Real per-pipeline blend + // modes land alongside the real pipelines in later phases. + let color_blend_attachment = vk::PipelineColorBlendAttachmentState::default() + .blend_enable(true) + .src_color_blend_factor(vk::BlendFactor::ONE) + .dst_color_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .color_blend_op(vk::BlendOp::ADD) + .src_alpha_blend_factor(vk::BlendFactor::ONE) + .dst_alpha_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .alpha_blend_op(vk::BlendOp::ADD) + .color_write_mask(vk::ColorComponentFlags::RGBA); + let color_blend_attachments = [color_blend_attachment]; + let color_blend = vk::PipelineColorBlendStateCreateInfo::default() + .attachments(&color_blend_attachments); + + let dynamic_states = [vk::DynamicState::VIEWPORT, vk::DynamicState::SCISSOR]; + let dynamic_state = + vk::PipelineDynamicStateCreateInfo::default().dynamic_states(&dynamic_states); + + // Dynamic rendering — no `VkRenderPass` needed. Just declare the + // color attachment format the pipeline will write to. + let color_attachment_formats = [color_format]; + let mut rendering = vk::PipelineRenderingCreateInfo::default() + .color_attachment_formats(&color_attachment_formats); + + let pipeline_info = vk::GraphicsPipelineCreateInfo::default() + .stages(&stages) + .vertex_input_state(&vertex_input) + .input_assembly_state(&input_assembly) + .viewport_state(&viewport_state) + .rasterization_state(&rasterization) + .multisample_state(&multisample) + .color_blend_state(&color_blend) + .dynamic_state(&dynamic_state) + .layout(layout) + .push_next(&mut rendering); + + unsafe { + device + .create_graphics_pipelines(pipeline_cache, &[pipeline_info], None) + .map_err(|(_, e)| e) + .expect("create_graphics_pipelines(bootstrap)")[0] + } +} + +// ======================================================================= +// Quad pipeline helpers +// ======================================================================= + +fn create_quad_descriptor_set_layout(device: &ash::Device) -> vk::DescriptorSetLayout { + let bindings = [vk::DescriptorSetLayoutBinding::default() + .binding(0) + .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::VERTEX | vk::ShaderStageFlags::FRAGMENT)]; + let info = vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings); + unsafe { + device + .create_descriptor_set_layout(&info, None) + .expect("create_descriptor_set_layout(quad)") + } +} + +fn create_quad_descriptor_pool(device: &ash::Device) -> vk::DescriptorPool { + let sizes = [vk::DescriptorPoolSize { + ty: vk::DescriptorType::UNIFORM_BUFFER, + descriptor_count: FRAMES_IN_FLIGHT as u32, + }]; + let info = vk::DescriptorPoolCreateInfo::default() + .max_sets(FRAMES_IN_FLIGHT as u32) + .pool_sizes(&sizes); + unsafe { + device + .create_descriptor_pool(&info, None) + .expect("create_descriptor_pool(quad)") + } +} + +fn allocate_quad_descriptor_sets( + device: &ash::Device, + pool: vk::DescriptorPool, + layout: vk::DescriptorSetLayout, +) -> [vk::DescriptorSet; FRAMES_IN_FLIGHT] { + let layouts = [layout; FRAMES_IN_FLIGHT]; + let info = vk::DescriptorSetAllocateInfo::default() + .descriptor_pool(pool) + .set_layouts(&layouts); + let sets = unsafe { + device + .allocate_descriptor_sets(&info) + .expect("allocate_descriptor_sets(quad)") + }; + let mut out = [vk::DescriptorSet::null(); FRAMES_IN_FLIGHT]; + out.copy_from_slice(&sets); + out +} + +fn update_quad_descriptor_set( + device: &ash::Device, + set: vk::DescriptorSet, + uniform: &VulkanBuffer, +) { + let uniform_info = vk::DescriptorBufferInfo::default() + .buffer(uniform.handle()) + .offset(0) + .range(uniform.size()); + let infos = [uniform_info]; + let write = vk::WriteDescriptorSet::default() + .dst_set(set) + .dst_binding(0) + .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER) + .buffer_info(&infos); + unsafe { + device.update_descriptor_sets(&[write], &[]); + } +} + +fn create_quad_pipeline_layout( + device: &ash::Device, + set_layout: vk::DescriptorSetLayout, +) -> vk::PipelineLayout { + let set_layouts = [set_layout]; + let info = vk::PipelineLayoutCreateInfo::default().set_layouts(&set_layouts); + unsafe { + device + .create_pipeline_layout(&info, None) + .expect("create_pipeline_layout(quad)") + } +} + +fn build_quad_pipeline( + device: &ash::Device, + pipeline_cache: vk::PipelineCache, + layout: vk::PipelineLayout, + color_format: vk::Format, +) -> vk::Pipeline { + let vert = create_shader_module(device, QUAD_VERT_SPV); + let frag = create_shader_module(device, QUAD_FRAG_SPV); + + let entry = c"main"; + let stages = [ + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::VERTEX) + .module(vert) + .name(entry), + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::FRAGMENT) + .module(frag) + .name(entry), + ]; + + // Vertex input mirrors `QuadInstance` (96 bytes, 8 attributes). + let bindings = [vk::VertexInputBindingDescription::default() + .binding(0) + .stride(std::mem::size_of::() as u32) + .input_rate(vk::VertexInputRate::INSTANCE)]; + let attrs = [ + // 0: pos vec3 @ 0 + vk::VertexInputAttributeDescription::default() + .location(0) + .binding(0) + .format(vk::Format::R32G32B32_SFLOAT) + .offset(0), + // 1: color vec4 @ 12 + vk::VertexInputAttributeDescription::default() + .location(1) + .binding(0) + .format(vk::Format::R32G32B32A32_SFLOAT) + .offset(12), + // 2: uv_rect vec4 @ 28 + vk::VertexInputAttributeDescription::default() + .location(2) + .binding(0) + .format(vk::Format::R32G32B32A32_SFLOAT) + .offset(28), + // 3: layers ivec2 @ 44 + vk::VertexInputAttributeDescription::default() + .location(3) + .binding(0) + .format(vk::Format::R32G32_SINT) + .offset(44), + // 4: size vec2 @ 52 + vk::VertexInputAttributeDescription::default() + .location(4) + .binding(0) + .format(vk::Format::R32G32_SFLOAT) + .offset(52), + // 5: corner_radii vec4 @ 60 + vk::VertexInputAttributeDescription::default() + .location(5) + .binding(0) + .format(vk::Format::R32G32B32A32_SFLOAT) + .offset(60), + // 6: underline_style i32 @ 76 + vk::VertexInputAttributeDescription::default() + .location(6) + .binding(0) + .format(vk::Format::R32_SINT) + .offset(76), + // 7: clip_rect vec4 @ 80 + vk::VertexInputAttributeDescription::default() + .location(7) + .binding(0) + .format(vk::Format::R32G32B32A32_SFLOAT) + .offset(80), + ]; + let vertex_input = vk::PipelineVertexInputStateCreateInfo::default() + .vertex_binding_descriptions(&bindings) + .vertex_attribute_descriptions(&attrs); + + let input_assembly = vk::PipelineInputAssemblyStateCreateInfo::default() + .topology(vk::PrimitiveTopology::TRIANGLE_STRIP); + let viewport_state = vk::PipelineViewportStateCreateInfo::default() + .viewport_count(1) + .scissor_count(1); + let rasterization = vk::PipelineRasterizationStateCreateInfo::default() + .polygon_mode(vk::PolygonMode::FILL) + .cull_mode(vk::CullModeFlags::NONE) + .front_face(vk::FrontFace::COUNTER_CLOCKWISE) + .line_width(1.0); + let multisample = vk::PipelineMultisampleStateCreateInfo::default() + .rasterization_samples(vk::SampleCountFlags::TYPE_1); + + // Same blend as Metal/wgpu rich-text pipeline: + // src_color: SrcAlpha, dst_color: OneMinusSrcAlpha (gamma-space + // alpha blending — matches sugarloaf's other pipelines). + let blend_attachment = vk::PipelineColorBlendAttachmentState::default() + .blend_enable(true) + .src_color_blend_factor(vk::BlendFactor::SRC_ALPHA) + .dst_color_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .color_blend_op(vk::BlendOp::ADD) + .src_alpha_blend_factor(vk::BlendFactor::ONE) + .dst_alpha_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .alpha_blend_op(vk::BlendOp::ADD) + .color_write_mask(vk::ColorComponentFlags::RGBA); + let blend_attachments = [blend_attachment]; + let color_blend = + vk::PipelineColorBlendStateCreateInfo::default().attachments(&blend_attachments); + + let dynamic_states = [vk::DynamicState::VIEWPORT, vk::DynamicState::SCISSOR]; + let dynamic_state = + vk::PipelineDynamicStateCreateInfo::default().dynamic_states(&dynamic_states); + + let color_attachment_formats = [color_format]; + let mut rendering = vk::PipelineRenderingCreateInfo::default() + .color_attachment_formats(&color_attachment_formats); + + let pipeline_info = vk::GraphicsPipelineCreateInfo::default() + .stages(&stages) + .vertex_input_state(&vertex_input) + .input_assembly_state(&input_assembly) + .viewport_state(&viewport_state) + .rasterization_state(&rasterization) + .multisample_state(&multisample) + .color_blend_state(&color_blend) + .dynamic_state(&dynamic_state) + .layout(layout) + .push_next(&mut rendering); + + let pipeline = unsafe { + device + .create_graphics_pipelines(pipeline_cache, &[pipeline_info], None) + .map_err(|(_, e)| e) + .expect("create_graphics_pipelines(quad)")[0] + }; + unsafe { + device.destroy_shader_module(vert, None); + device.destroy_shader_module(frag, None); + } + pipeline +} + +// ======================================================================= +// Image pipeline helpers +// ======================================================================= + +fn create_image_uniform_descriptor_set_layout( + device: &ash::Device, +) -> vk::DescriptorSetLayout { + // Same shape as the quad pipeline's uniform set: one + // UNIFORM_BUFFER at binding 0, visible to vertex + fragment. + let bindings = [vk::DescriptorSetLayoutBinding::default() + .binding(0) + .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::VERTEX | vk::ShaderStageFlags::FRAGMENT)]; + let info = vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings); + unsafe { + device + .create_descriptor_set_layout(&info, None) + .expect("create_descriptor_set_layout(image uniform)") + } +} + +fn create_image_texture_descriptor_set_layout( + device: &ash::Device, +) -> vk::DescriptorSetLayout { + let bindings = [vk::DescriptorSetLayoutBinding::default() + .binding(0) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::FRAGMENT)]; + let info = vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings); + unsafe { + device + .create_descriptor_set_layout(&info, None) + .expect("create_descriptor_set_layout(image texture)") + } +} + +fn create_image_uniform_descriptor_pool(device: &ash::Device) -> vk::DescriptorPool { + let sizes = [vk::DescriptorPoolSize { + ty: vk::DescriptorType::UNIFORM_BUFFER, + descriptor_count: FRAMES_IN_FLIGHT as u32, + }]; + let info = vk::DescriptorPoolCreateInfo::default() + .max_sets(FRAMES_IN_FLIGHT as u32) + .pool_sizes(&sizes); + unsafe { + device + .create_descriptor_pool(&info, None) + .expect("create_descriptor_pool(image uniform)") + } +} + +fn create_image_pipeline_layout( + device: &ash::Device, + uniform_layout: vk::DescriptorSetLayout, + texture_layout: vk::DescriptorSetLayout, +) -> vk::PipelineLayout { + let set_layouts = [uniform_layout, texture_layout]; + let info = vk::PipelineLayoutCreateInfo::default().set_layouts(&set_layouts); + unsafe { + device + .create_pipeline_layout(&info, None) + .expect("create_pipeline_layout(image)") + } +} + +fn create_image_sampler(device: &ash::Device) -> vk::Sampler { + // Linear filtering for smooth scaling of background images and + // kitty graphics. ClampToEdge addressing prevents bleeding at + // the texture edges. + let info = vk::SamplerCreateInfo::default() + .mag_filter(vk::Filter::LINEAR) + .min_filter(vk::Filter::LINEAR) + .mipmap_mode(vk::SamplerMipmapMode::LINEAR) + .address_mode_u(vk::SamplerAddressMode::CLAMP_TO_EDGE) + .address_mode_v(vk::SamplerAddressMode::CLAMP_TO_EDGE) + .address_mode_w(vk::SamplerAddressMode::CLAMP_TO_EDGE); + unsafe { + device + .create_sampler(&info, None) + .expect("create_sampler(image)") + } +} + +fn build_image_pipeline( + device: &ash::Device, + pipeline_cache: vk::PipelineCache, + layout: vk::PipelineLayout, + color_format: vk::Format, +) -> vk::Pipeline { + let vert = create_shader_module(device, IMAGE_VERT_SPV); + let frag = create_shader_module(device, IMAGE_FRAG_SPV); + + let entry = c"main"; + let stages = [ + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::VERTEX) + .module(vert) + .name(entry), + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::FRAGMENT) + .module(frag) + .name(entry), + ]; + + let bindings = [vk::VertexInputBindingDescription::default() + .binding(0) + .stride(std::mem::size_of::() as u32) + .input_rate(vk::VertexInputRate::INSTANCE)]; + let attrs = [ + // 0: dest_pos vec2 @ 0 + vk::VertexInputAttributeDescription::default() + .location(0) + .binding(0) + .format(vk::Format::R32G32_SFLOAT) + .offset(0), + // 1: dest_size vec2 @ 8 + vk::VertexInputAttributeDescription::default() + .location(1) + .binding(0) + .format(vk::Format::R32G32_SFLOAT) + .offset(8), + // 2: source_rect vec4 @ 16 + vk::VertexInputAttributeDescription::default() + .location(2) + .binding(0) + .format(vk::Format::R32G32B32A32_SFLOAT) + .offset(16), + ]; + let vertex_input = vk::PipelineVertexInputStateCreateInfo::default() + .vertex_binding_descriptions(&bindings) + .vertex_attribute_descriptions(&attrs); + + let input_assembly = vk::PipelineInputAssemblyStateCreateInfo::default() + .topology(vk::PrimitiveTopology::TRIANGLE_STRIP); + let viewport_state = vk::PipelineViewportStateCreateInfo::default() + .viewport_count(1) + .scissor_count(1); + let rasterization = vk::PipelineRasterizationStateCreateInfo::default() + .polygon_mode(vk::PolygonMode::FILL) + .cull_mode(vk::CullModeFlags::NONE) + .front_face(vk::FrontFace::COUNTER_CLOCKWISE) + .line_width(1.0); + let multisample = vk::PipelineMultisampleStateCreateInfo::default() + .rasterization_samples(vk::SampleCountFlags::TYPE_1); + + // Image fragment returns premultiplied RGBA — source factor ONE. + let blend_attachment = vk::PipelineColorBlendAttachmentState::default() + .blend_enable(true) + .src_color_blend_factor(vk::BlendFactor::ONE) + .dst_color_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .color_blend_op(vk::BlendOp::ADD) + .src_alpha_blend_factor(vk::BlendFactor::ONE) + .dst_alpha_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .alpha_blend_op(vk::BlendOp::ADD) + .color_write_mask(vk::ColorComponentFlags::RGBA); + let blend_attachments = [blend_attachment]; + let color_blend = + vk::PipelineColorBlendStateCreateInfo::default().attachments(&blend_attachments); + + let dynamic_states = [vk::DynamicState::VIEWPORT, vk::DynamicState::SCISSOR]; + let dynamic_state = + vk::PipelineDynamicStateCreateInfo::default().dynamic_states(&dynamic_states); + + let color_attachment_formats = [color_format]; + let mut rendering = vk::PipelineRenderingCreateInfo::default() + .color_attachment_formats(&color_attachment_formats); + + let pipeline_info = vk::GraphicsPipelineCreateInfo::default() + .stages(&stages) + .vertex_input_state(&vertex_input) + .input_assembly_state(&input_assembly) + .viewport_state(&viewport_state) + .rasterization_state(&rasterization) + .multisample_state(&multisample) + .color_blend_state(&color_blend) + .dynamic_state(&dynamic_state) + .layout(layout) + .push_next(&mut rendering); + + let pipeline = unsafe { + device + .create_graphics_pipelines(pipeline_cache, &[pipeline_info], None) + .map_err(|(_, e)| e) + .expect("create_graphics_pipelines(image)")[0] + }; + unsafe { + device.destroy_shader_module(vert, None); + device.destroy_shader_module(frag, None); + } + pipeline +} + +// ======================================================================= +// Per-image texture: device-local image + per-image descriptor set +// ======================================================================= + +/// One uploaded image (background, kitty graphic, sixel) — owns its +/// `vk::Image`, view, memory, descriptor pool, and descriptor set. +/// The descriptor set's binding 0 is wired to the image's view + +/// the `VulkanRenderer`'s shared sampler at construction time, so +/// the renderer's `render_*` paths can bind it directly without +/// touching descriptor pools per draw. +pub struct VulkanImageTexture { + pub image: VulkanImage, + descriptor_pool: vk::DescriptorPool, + pub descriptor_set: vk::DescriptorSet, + device: ash::Device, +} + +impl VulkanImageTexture { + /// Synchronously upload `pixels` (RGBA8) into a fresh device-local + /// image, transition to `SHADER_READ_ONLY_OPTIMAL`, and create a + /// descriptor set bound to (image_view, sampler). + /// + /// The submit-and-wait is the right choice for one-shot uploads + /// (background image, set once at config-load). Per-frame upload + /// paths (kitty graphics) want a different code path that + /// piggy-backs on the per-frame command buffer. + pub fn upload_rgba( + ctx: &VulkanContext, + pixels: &[u8], + width: u32, + height: u32, + descriptor_set_layout: vk::DescriptorSetLayout, + sampler: vk::Sampler, + ) -> Self { + let device = ctx.device().clone(); + + // `R8G8B8A8_SRGB` (vs `R8G8B8A8_UNORM`) tells the GPU to + // sRGB-decode bytes at sample time. With bilinear filtering + // enabled on the sampler this means interpolation happens in + // *linear* light — without it, scaled image edges come out + // visibly dark (gamma-space midtones), matching the same + // choice on `RGBA8Unorm_sRGB` in `context::metal`. + let image = ctx.allocate_sampled_image( + width, + height, + vk::Format::R8G8B8A8_SRGB, + vk::ImageUsageFlags::TRANSFER_DST | vk::ImageUsageFlags::SAMPLED, + ); + + // Staging buffer. + let staging_size = (width as usize) * (height as usize) * 4; + let staging = ctx.allocate_host_visible_buffer( + staging_size as u64, + vk::BufferUsageFlags::TRANSFER_SRC, + ); + unsafe { + std::ptr::copy_nonoverlapping( + pixels.as_ptr(), + staging.as_mut_ptr(), + staging_size, + ); + } + + // One-shot transfer: barrier → copy → barrier. + let img_handle = image.handle(); + let staging_handle = staging.handle(); + ctx.submit_oneshot(|cmd| unsafe { + let to_dst = vk::ImageMemoryBarrier2::default() + .src_stage_mask(vk::PipelineStageFlags2::TOP_OF_PIPE) + .src_access_mask(vk::AccessFlags2::empty()) + .dst_stage_mask(vk::PipelineStageFlags2::COPY) + .dst_access_mask(vk::AccessFlags2::TRANSFER_WRITE) + .old_layout(vk::ImageLayout::UNDEFINED) + .new_layout(vk::ImageLayout::TRANSFER_DST_OPTIMAL) + .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .image(img_handle) + .subresource_range( + vk::ImageSubresourceRange::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .base_mip_level(0) + .level_count(1) + .base_array_layer(0) + .layer_count(1), + ); + let barriers = [to_dst]; + let dep = vk::DependencyInfo::default().image_memory_barriers(&barriers); + device.cmd_pipeline_barrier2(cmd, &dep); + + let region = vk::BufferImageCopy::default() + .image_subresource( + vk::ImageSubresourceLayers::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .mip_level(0) + .base_array_layer(0) + .layer_count(1), + ) + .image_extent(vk::Extent3D { + width, + height, + depth: 1, + }); + device.cmd_copy_buffer_to_image( + cmd, + staging_handle, + img_handle, + vk::ImageLayout::TRANSFER_DST_OPTIMAL, + &[region], + ); + + let to_read = vk::ImageMemoryBarrier2::default() + .src_stage_mask(vk::PipelineStageFlags2::COPY) + .src_access_mask(vk::AccessFlags2::TRANSFER_WRITE) + .dst_stage_mask(vk::PipelineStageFlags2::FRAGMENT_SHADER) + .dst_access_mask(vk::AccessFlags2::SHADER_READ) + .old_layout(vk::ImageLayout::TRANSFER_DST_OPTIMAL) + .new_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL) + .src_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .dst_queue_family_index(vk::QUEUE_FAMILY_IGNORED) + .image(img_handle) + .subresource_range( + vk::ImageSubresourceRange::default() + .aspect_mask(vk::ImageAspectFlags::COLOR) + .base_mip_level(0) + .level_count(1) + .base_array_layer(0) + .layer_count(1), + ); + let barriers = [to_read]; + let dep = vk::DependencyInfo::default().image_memory_barriers(&barriers); + device.cmd_pipeline_barrier2(cmd, &dep); + }); + // Staging buffer drops here — submit_oneshot already waited. + + // Per-image descriptor pool + set. + let pool_sizes = [vk::DescriptorPoolSize { + ty: vk::DescriptorType::COMBINED_IMAGE_SAMPLER, + descriptor_count: 1, + }]; + let pool_info = vk::DescriptorPoolCreateInfo::default() + .max_sets(1) + .pool_sizes(&pool_sizes); + let descriptor_pool = unsafe { + device + .create_descriptor_pool(&pool_info, None) + .expect("create_descriptor_pool(image texture)") + }; + let layouts = [descriptor_set_layout]; + let alloc_info = vk::DescriptorSetAllocateInfo::default() + .descriptor_pool(descriptor_pool) + .set_layouts(&layouts); + let descriptor_set = unsafe { + device + .allocate_descriptor_sets(&alloc_info) + .expect("allocate_descriptor_sets(image texture)")[0] + }; + + let image_info = vk::DescriptorImageInfo::default() + .sampler(sampler) + .image_view(image.view()) + .image_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL); + let infos = [image_info]; + let write = vk::WriteDescriptorSet::default() + .dst_set(descriptor_set) + .dst_binding(0) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .image_info(&infos); + unsafe { + device.update_descriptor_sets(&[write], &[]); + } + + Self { + image, + descriptor_pool, + descriptor_set, + device, + } + } +} + +impl Drop for VulkanImageTexture { + fn drop(&mut self) { + unsafe { + // Pool destruction frees the descriptor set; image drops + // itself. + self.device + .destroy_descriptor_pool(self.descriptor_pool, None); + } + } +} + +// ======================================================================= +// Geometry pipeline (per-vertex `Vertex` for non-quad draws) +// ======================================================================= + +fn build_geometry_pipeline( + device: &ash::Device, + pipeline_cache: vk::PipelineCache, + layout: vk::PipelineLayout, + color_format: vk::Format, +) -> vk::Pipeline { + // Reuses the QUAD fragment shader — Vertex output structure is + // intentionally identical to QuadInstance's vertex output. + let vert = create_shader_module(device, GEOMETRY_VERT_SPV); + let frag = create_shader_module(device, QUAD_FRAG_SPV); + + let entry = c"main"; + let stages = [ + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::VERTEX) + .module(vert) + .name(entry), + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::FRAGMENT) + .module(frag) + .name(entry), + ]; + + // Per-vertex (NOT instanced) — matches `Vertex` (88 bytes). + let bindings = [vk::VertexInputBindingDescription::default() + .binding(0) + .stride(std::mem::size_of::() as u32) + .input_rate(vk::VertexInputRate::VERTEX)]; + let attrs = [ + // 0: pos vec3 @ 0 + vk::VertexInputAttributeDescription::default() + .location(0) + .binding(0) + .format(vk::Format::R32G32B32_SFLOAT) + .offset(0), + // 1: color vec4 @ 12 + vk::VertexInputAttributeDescription::default() + .location(1) + .binding(0) + .format(vk::Format::R32G32B32A32_SFLOAT) + .offset(12), + // 2: uv vec2 @ 28 + vk::VertexInputAttributeDescription::default() + .location(2) + .binding(0) + .format(vk::Format::R32G32_SFLOAT) + .offset(28), + // 3: layers ivec2 @ 36 + vk::VertexInputAttributeDescription::default() + .location(3) + .binding(0) + .format(vk::Format::R32G32_SINT) + .offset(36), + // 4: corner_radii vec4 @ 44 + vk::VertexInputAttributeDescription::default() + .location(4) + .binding(0) + .format(vk::Format::R32G32B32A32_SFLOAT) + .offset(44), + // 5: rect_size vec2 @ 60 + vk::VertexInputAttributeDescription::default() + .location(5) + .binding(0) + .format(vk::Format::R32G32_SFLOAT) + .offset(60), + // 6: underline_style i32 @ 68 + vk::VertexInputAttributeDescription::default() + .location(6) + .binding(0) + .format(vk::Format::R32_SINT) + .offset(68), + // 7: clip_rect vec4 @ 72 + vk::VertexInputAttributeDescription::default() + .location(7) + .binding(0) + .format(vk::Format::R32G32B32A32_SFLOAT) + .offset(72), + ]; + let vertex_input = vk::PipelineVertexInputStateCreateInfo::default() + .vertex_binding_descriptions(&bindings) + .vertex_attribute_descriptions(&attrs); + + // TRIANGLE_LIST — emit path tessellates polygons / arcs / lines + // into independent triangles (3 vertices each, no strip). + let input_assembly = vk::PipelineInputAssemblyStateCreateInfo::default() + .topology(vk::PrimitiveTopology::TRIANGLE_LIST); + let viewport_state = vk::PipelineViewportStateCreateInfo::default() + .viewport_count(1) + .scissor_count(1); + let rasterization = vk::PipelineRasterizationStateCreateInfo::default() + .polygon_mode(vk::PolygonMode::FILL) + .cull_mode(vk::CullModeFlags::NONE) + .front_face(vk::FrontFace::COUNTER_CLOCKWISE) + .line_width(1.0); + let multisample = vk::PipelineMultisampleStateCreateInfo::default() + .rasterization_samples(vk::SampleCountFlags::TYPE_1); + + // Same blend as the quad pipeline — gamma-space SrcAlpha / + // OneMinusSrcAlpha, matching every other sugarloaf pipeline. + let blend_attachment = vk::PipelineColorBlendAttachmentState::default() + .blend_enable(true) + .src_color_blend_factor(vk::BlendFactor::SRC_ALPHA) + .dst_color_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .color_blend_op(vk::BlendOp::ADD) + .src_alpha_blend_factor(vk::BlendFactor::ONE) + .dst_alpha_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .alpha_blend_op(vk::BlendOp::ADD) + .color_write_mask(vk::ColorComponentFlags::RGBA); + let blend_attachments = [blend_attachment]; + let color_blend = + vk::PipelineColorBlendStateCreateInfo::default().attachments(&blend_attachments); + + let dynamic_states = [vk::DynamicState::VIEWPORT, vk::DynamicState::SCISSOR]; + let dynamic_state = + vk::PipelineDynamicStateCreateInfo::default().dynamic_states(&dynamic_states); + + let color_attachment_formats = [color_format]; + let mut rendering = vk::PipelineRenderingCreateInfo::default() + .color_attachment_formats(&color_attachment_formats); + + let pipeline_info = vk::GraphicsPipelineCreateInfo::default() + .stages(&stages) + .vertex_input_state(&vertex_input) + .input_assembly_state(&input_assembly) + .viewport_state(&viewport_state) + .rasterization_state(&rasterization) + .multisample_state(&multisample) + .color_blend_state(&color_blend) + .dynamic_state(&dynamic_state) + .layout(layout) + .push_next(&mut rendering); + + let pipeline = unsafe { + device + .create_graphics_pipelines(pipeline_cache, &[pipeline_info], None) + .map_err(|(_, e)| e) + .expect("create_graphics_pipelines(geometry)")[0] + }; + unsafe { + device.destroy_shader_module(vert, None); + device.destroy_shader_module(frag, None); + } + pipeline +} diff --git a/sugarloaf/src/sugarloaf.rs b/sugarloaf/src/sugarloaf.rs index bda808fa..dc2399bb 100644 --- a/sugarloaf/src/sugarloaf.rs +++ b/sugarloaf/src/sugarloaf.rs @@ -3,6 +3,7 @@ pub mod primitives; pub mod state; use crate::components::core::image::Handle; +#[cfg(feature = "wgpu")] use crate::components::filters::{Filter, FiltersBrush}; use crate::font::{fonts::SugarloafFont, FontLibrary}; use crate::font_cache::{compute_advance, resolve_with, FontCache, ResolvedGlyph}; @@ -22,7 +23,11 @@ use raw_window_handle::{ use state::SugarState; pub struct Sugarloaf<'a> { - pub ctx: Context<'a>, + // NOTE: field order is load-bearing for the Vulkan backend. Rust + // drops struct fields in declaration order, and any field that + // owns Vulkan handles (renderer, text, eventually grids) must drop + // BEFORE `ctx` — destroying handles after the device has been + // torn down would crash the driver. `ctx` is therefore last. renderer: Renderer, state: state::SugarState, /// Input colorspace configured at construction. Exposed via @@ -30,9 +35,10 @@ pub struct Sugarloaf<'a> { /// value into its own uniform — keeps grid + quad-fill pipelines /// producing byte-identical framebuffer colors. colorspace: Colorspace, - pub background_color: Option, + pub background_color: Option, pub background_image: Option, pub graphics: Graphics, + #[cfg(feature = "wgpu")] filters_brush: Option, /// Pixel data for standalone image textures, keyed by ImageId. pub image_data: rustc_hash::FxHashMap, @@ -55,6 +61,10 @@ pub struct Sugarloaf<'a> { /// graphics frontend path; read by the renderer's image pass. pub image_overlays: rustc_hash::FxHashMap>, + /// Owned context (device + swapchain + queue). Last so the device + /// outlives every Vulkan-handle-owning field above. See note at + /// the top of the struct. + pub ctx: Context<'a>, } #[derive(Debug)] @@ -87,15 +97,59 @@ pub struct SugarloafWindow { } pub enum SugarloafBackend { + /// `wgpu` umbrella backend (Vulkan / Metal / DX12 / GL / WebGPU, + /// whichever wgpu picks). Only available when sugarloaf is built + /// with the `wgpu` feature; on Linux/macOS the native `Vulkan` + /// and `Metal` variants are preferred and wgpu can be omitted + /// entirely from the dep tree. Always available on Windows and + /// WASM (the native backends don't cover those yet). + #[cfg(feature = "wgpu")] Wgpu(wgpu::Backends), #[cfg(target_os = "macos")] Metal, + /// Native Vulkan via `ash`. Linux only for now (Windows would need a + /// `khr::win32_surface` branch in `context::vulkan::create_surface`). + /// Mirrors the Metal backend in scope: no librashader filters. + #[cfg(target_os = "linux")] + Vulkan, /// CPU rendering via tiny-skia + softbuffer. Cpu, } +/// RGBA color in linear-light 0..1 space. Mirrors `wgpu::Color`'s +/// shape so callers don't have to depend on `wgpu`. Sugarloaf's +/// public API takes/returns this type; the wgpu render path +/// converts at the boundary. +#[repr(C)] +#[derive(Clone, Copy, Debug, PartialEq)] +pub struct Color { + pub r: f64, + pub g: f64, + pub b: f64, + pub a: f64, +} + +impl Color { + pub const TRANSPARENT: Self = Self { r: 0.0, g: 0.0, b: 0.0, a: 0.0 }; + pub const BLACK: Self = Self { r: 0.0, g: 0.0, b: 0.0, a: 1.0 }; + pub const WHITE: Self = Self { r: 1.0, g: 1.0, b: 1.0, a: 1.0 }; +} + +#[cfg(feature = "wgpu")] +impl From for wgpu::Color { + fn from(c: Color) -> Self { + wgpu::Color { r: c.r, g: c.g, b: c.b, a: c.a } + } +} + +#[cfg(feature = "wgpu")] +impl From for Color { + fn from(c: wgpu::Color) -> Self { + Color { r: c.r, g: c.g, b: c.b, a: c.a } + } +} + pub struct SugarloafRenderer { - pub power_preference: wgpu::PowerPreference, pub backend: SugarloafBackend, pub font_features: Option>, pub colorspace: Colorspace, @@ -121,18 +175,38 @@ impl Default for Colorspace { impl Default for SugarloafRenderer { fn default() -> SugarloafRenderer { - #[cfg(target_arch = "wasm32")] + #[cfg(all(target_arch = "wasm32", feature = "wgpu"))] let default_backend = SugarloafBackend::Wgpu(wgpu::Backends::BROWSER_WEBGPU | wgpu::Backends::GL); - #[cfg(not(any(target_arch = "wasm32", target_os = "macos")))] + // Linux defaults to the native Vulkan backend (ash). Mirrors the + // macOS default to native Metal — both skip the wgpu translation + // layer and the librashader filter chain. Other non-macOS desktop + // targets (Windows, BSDs) need the `wgpu` feature enabled and + // fall back to the wgpu umbrella backend. + #[cfg(all(target_os = "linux", not(target_arch = "wasm32")))] + let default_backend = SugarloafBackend::Vulkan; + + #[cfg(all( + not(target_arch = "wasm32"), + not(target_os = "macos"), + not(target_os = "linux"), + feature = "wgpu", + ))] let default_backend = SugarloafBackend::Wgpu(wgpu::Backends::all()); + #[cfg(all( + not(target_arch = "wasm32"), + not(target_os = "macos"), + not(target_os = "linux"), + not(feature = "wgpu"), + ))] + let default_backend = SugarloafBackend::Cpu; + #[cfg(all(target_os = "macos", not(target_arch = "wasm32")))] let default_backend = SugarloafBackend::Metal; SugarloafRenderer { - power_preference: wgpu::PowerPreference::HighPerformance, backend: default_backend, font_features: None, colorspace: Colorspace::default(), @@ -190,10 +264,11 @@ impl Sugarloaf<'_> { state, ctx, colorspace, - background_color: Some(wgpu::Color::BLACK), + background_color: Some(Color::BLACK), background_image: None, renderer, graphics: Graphics::default(), + #[cfg(feature = "wgpu")] filters_brush: None, image_data: rustc_hash::FxHashMap::default(), cpu_cache: crate::renderer::cpu::CpuCache::new(), @@ -394,6 +469,11 @@ impl Sugarloaf<'_> { } #[inline] + /// Install librashader CRT/scanline filters. Only available with + /// the `wgpu` feature — librashader's runtime is wgpu-only + /// upstream, and the native Metal/Vulkan backends ignore filter + /// requests entirely. + #[cfg(feature = "wgpu")] pub fn update_filters(&mut self, filters: &[Filter]) { if filters.is_empty() { self.filters_brush = None; @@ -410,7 +490,7 @@ impl Sugarloaf<'_> { } #[inline] - pub fn set_background_color(&mut self, color: Option) -> &mut Self { + pub fn set_background_color(&mut self, color: Option) -> &mut Self { self.background_color = color; self } @@ -1035,6 +1115,7 @@ impl Sugarloaf<'_> { ); match self.ctx.inner { + #[cfg(feature = "wgpu")] crate::context::ContextType::Wgpu(_) => { self.render_wgpu(grids); } @@ -1042,9 +1123,15 @@ impl Sugarloaf<'_> { crate::context::ContextType::Metal(_) => { self.render_metal(grids); } + #[cfg(target_os = "linux")] + crate::context::ContextType::Vulkan(_) => { + self.render_vulkan(grids); + } crate::context::ContextType::Cpu(_) => { self.render_cpu(); } + #[cfg(not(feature = "wgpu"))] + crate::context::ContextType::_Phantom(_) => unreachable!(), } } @@ -1090,7 +1177,130 @@ impl Sugarloaf<'_> { self.reset(); } + /// Drive a native Vulkan frame. Drives the entire frame from the + /// outside so grid passes, the optional bootstrap rect, and (later) + /// rich-text / images / UI text overlays all share one + /// dynamic-rendering pass and one swapchain present. + /// + /// Order of operations inside the pass: + /// 1. swapchain image barrier `UNDEFINED → COLOR_ATTACHMENT_OPTIMAL` + /// 2. `cmd_begin_rendering` with `LOAD_OP_CLEAR(bg)` + /// 3. set viewport + scissor + /// 4. per-panel `GridRenderer::render_vulkan` (Phase 3+) + /// 5. `VulkanRenderer::draw_bootstrap` (debug rect, gated by env) + /// 6. `cmd_end_rendering` + /// 7. swapchain image barrier `COLOR_ATTACHMENT_OPTIMAL → PRESENT_SRC_KHR` + /// 8. `present_frame` (queue submit + present) + #[inline] + #[cfg(target_os = "linux")] + pub fn render_vulkan( + &mut self, + grids: &mut [(&mut crate::grid::GridRenderer, crate::grid::GridUniforms)], + ) { + use crate::renderer::vulkan as vkr; + + let ctx = match &mut self.ctx.inner { + crate::context::ContextType::Vulkan(v) => v, + _ => return, + }; + + let frame = match ctx.acquire_frame() { + Some(f) => f, + None => { + // Swapchain was recreated this turn — skip drawing, + // try again next frame. + self.reset(); + return; + } + }; + + let bg = self + .background_color + .map(|c| [c.r as f32, c.g as f32, c.b as f32, c.a as f32]) + .unwrap_or([0.0, 0.0, 0.0, 1.0]); + + let device = ctx.device().clone(); + let cmd = frame.cmd_buffer; + + // Pre-pass: per-panel grid atlas uploads + UI text overlay + // atlas uploads. MUST happen before `cmd_begin_rendering` — + // `vkCmdCopyBufferToImage` is invalid inside a render pass. + self.text.init_vulkan(ctx); + for (grid, _) in grids.iter_mut() { + grid.prepare_vulkan(ctx, cmd, frame.slot); + } + self.text.prepare_vulkan(ctx, cmd, frame.slot); + + vkr::cmd_acquire_image_for_rendering(&device, cmd, frame.image); + + let color_attachment = vkr::build_color_attachment(&frame, bg); + let color_attachments = [color_attachment]; + let rendering_info = vkr::build_rendering_info(&frame, &color_attachments); + + // Negative viewport height: Vulkan's NDC has y pointing DOWN + // (-1 = top, +1 = bottom), but `orthographic_projection` is + // written for the OpenGL/Metal convention (y up). Without the + // flip, the text vertex shader would map pixel-y=0 to NDC y=+1 + // (bottom in Vulkan), and the whole grid would render + // upside-down with row 0 at the bottom. Negative height tells + // the rasterizer to invert y, making the matrix work + // unchanged. Requires Vulkan 1.1+ (core, no extension). + // + // The bg pipeline is unaffected — its fragment shader uses + // `gl_FragCoord`, which is always in framebuffer pixel + // coordinates with top-left origin regardless of viewport. + let viewport = ash::vk::Viewport { + x: 0.0, + y: frame.extent.height as f32, + width: frame.extent.width as f32, + height: -(frame.extent.height as f32), + min_depth: 0.0, + max_depth: 1.0, + }; + let scissor = ash::vk::Rect2D { + offset: ash::vk::Offset2D { x: 0, y: 0 }, + extent: frame.extent, + }; + unsafe { + device.cmd_begin_rendering(cmd, &rendering_info); + device.cmd_set_viewport(cmd, 0, &[viewport]); + device.cmd_set_scissor(cmd, 0, &[scissor]); + } + + // Per-panel grid passes — draw cell backgrounds + grid text + // underneath everything else. + for (grid, uniforms) in grids.iter_mut() { + grid.render_vulkan(ctx, cmd, frame.slot, uniforms); + } + + // Rich-text quad pass — `Sugarloaf::quad()` / `rect()` calls + // (command palette background, search overlay panel, + // assistant frame, etc.). Drawn AFTER grid so panel chrome + // covers cells underneath, BEFORE the text overlay so labels + // sit on top. + self.renderer.render_vulkan(cmd, &frame); + + // UI text overlay (tab titles, search overlay labels, + // command palette items, etc.). Drawn last so labels sit on + // top of the panel chrome. + self.text.render_vulkan( + cmd, + frame.slot, + [frame.extent.width as f32, frame.extent.height as f32], + ); + + unsafe { + device.cmd_end_rendering(cmd); + } + vkr::cmd_release_image_to_present(&device, cmd, frame.image); + + ctx.present_frame(frame); + + self.reset(); + } + #[inline] + #[cfg(feature = "wgpu")] pub fn render_wgpu( &mut self, grids: &mut [(&mut crate::grid::GridRenderer, crate::grid::GridUniforms)], @@ -1114,7 +1324,7 @@ impl Sugarloaf<'_> { { let load = if let Some(background_color) = self.background_color { - wgpu::LoadOp::Clear(background_color) + wgpu::LoadOp::Clear(background_color.into()) } else { wgpu::LoadOp::Load }; @@ -1156,10 +1366,8 @@ impl Sugarloaf<'_> { #[cfg(not(target_os = "macos"))] { self.text.init_wgpu(&ctx.device, &ctx.queue, ctx.format); - self.text.render_wgpu( - &mut rpass, - [ctx.size.width as f32, ctx.size.height as f32], - ); + self.text + .render_wgpu(&mut rpass, [ctx.size.width, ctx.size.height]); } } diff --git a/sugarloaf/src/text.rs b/sugarloaf/src/text.rs index f6082615..cd9ac9d2 100644 --- a/sugarloaf/src/text.rs +++ b/sugarloaf/src/text.rs @@ -139,6 +139,38 @@ struct TextWgpuState { instance_capacity: usize, } +#[cfg(target_os = "linux")] +struct TextVulkanState { + device: ash::Device, + instance: ash::Instance, + physical_device: ash::vk::PhysicalDevice, + /// Independent atlases owned by the UI text overlay — separate + /// from each `VulkanGridRenderer`'s atlases so an overlay glyph + /// doesn't have to compete for grid atlas space (and vice versa). + /// Same shape as the macOS / wgpu paths. + atlas_grayscale: crate::grid::vulkan::VulkanGlyphAtlas, + atlas_color: crate::grid::vulkan::VulkanGlyphAtlas, + /// Single sampler shared by both atlases. We only `texelFetch`, + /// so the sampler's filter / address mode don't matter — it just + /// has to exist for the COMBINED_IMAGE_SAMPLER descriptor. + sampler: ash::vk::Sampler, + uniform_buffers: + [crate::context::vulkan::VulkanBuffer; crate::context::vulkan::FRAMES_IN_FLIGHT], + /// Per-slot ring of instance buffers. Grow on demand, never shrink. + instance_buffers: [Option; + crate::context::vulkan::FRAMES_IN_FLIGHT], + instance_capacity: [usize; crate::context::vulkan::FRAMES_IN_FLIGHT], + descriptor_pool: ash::vk::DescriptorPool, + uniform_descriptor_set_layout: ash::vk::DescriptorSetLayout, + atlas_descriptor_set_layout: ash::vk::DescriptorSetLayout, + uniform_descriptor_sets: + [ash::vk::DescriptorSet; crate::context::vulkan::FRAMES_IN_FLIGHT], + atlas_descriptor_set: ash::vk::DescriptorSet, + + pipeline_layout: ash::vk::PipelineLayout, + pipeline: ash::vk::Pipeline, +} + // Text — the immediate-mode recorder owned by Sugarloaf pub struct Text { @@ -188,6 +220,8 @@ pub struct Text { font_data_cache: FxHashMap, #[cfg(not(target_os = "macos"))] wgpu: Option, + #[cfg(target_os = "linux")] + vulkan: Option, } impl Text { @@ -212,6 +246,8 @@ impl Text { font_data_cache: FxHashMap::default(), #[cfg(not(target_os = "macos"))] wgpu: None, + #[cfg(target_os = "linux")] + vulkan: None, } } @@ -543,9 +579,61 @@ impl Text { )) } - // ---- non-macOS (swash → WgpuGlyphAtlas) ---- + // ---- non-macOS (swash → VulkanGlyphAtlas or WgpuGlyphAtlas) ---- #[cfg(not(target_os = "macos"))] { + // Look up the slot first — backend-agnostic — and + // rasterize+insert into whichever atlas is initialized. + // Vulkan takes precedence on Linux when the Vulkan + // backend is active; wgpu is the fallback. + #[cfg(target_os = "linux")] + if self.vulkan.is_some() { + let state = self.vulkan.as_mut()?; + if let Some(s) = state.atlas_grayscale.lookup(key) { + return Some((s.x, s.y, s.w, s.h, s.bearing_x, s.bearing_y, false)); + } + if let Some(s) = state.atlas_color.lookup(key) { + return Some((s.x, s.y, s.w, s.h, s.bearing_x, s.bearing_y, true)); + } + + let font_entry = self.font_data_cache.get(&run.font_id)?.clone(); + let raw = rasterize_swash_glyph( + &mut self.scale_ctx, + &font_entry, + glyph_id, + run.size_u16 as f32, + run.synthetic_bold, + run.synthetic_italic, + self.font_library.inner.read().hinting, + )?; + let is_color = raw.is_color; + let raster = crate::grid::RasterizedGlyph { + width: raw.width.min(u16::MAX as u32) as u16, + height: raw.height.min(u16::MAX as u32) as u16, + bearing_x: raw.left.clamp(i16::MIN as i32, i16::MAX as i32) as i16, + bearing_y: { + let top_i16 = + raw.top.clamp(i16::MIN as i32, i16::MAX as i32) as i16; + run.ascent_px.saturating_sub(top_i16) + }, + bytes: &raw.bytes, + }; + let slot = if is_color { + state.atlas_color.insert(key, raster)? + } else { + state.atlas_grayscale.insert(key, raster)? + }; + return Some(( + slot.x, + slot.y, + slot.w, + slot.h, + slot.bearing_x, + slot.bearing_y, + is_color, + )); + } + let state = self.wgpu.as_mut()?; if let Some(s) = state.atlas_grayscale.lookup(key) { @@ -782,6 +870,120 @@ impl Text { render_pass.set_vertex_buffer(0, state.instance_buffer.slice(..)); render_pass.draw(0..4, 0..instance_count as u32); } + + // Vulkan GPU backend + + #[cfg(target_os = "linux")] + pub fn init_vulkan(&mut self, ctx: &crate::context::vulkan::VulkanContext) { + if self.vulkan.is_some() { + return; + } + let state = build_text_vulkan_state(ctx); + self.vulkan = Some(state); + } + + /// Pre-pass hook: drain pending atlas uploads into `cmd`. MUST be + /// called BEFORE `Sugarloaf::render_vulkan` opens its + /// dynamic-rendering pass (matches `GridRenderer::prepare_vulkan`). + /// No-op when neither atlas has pending uploads. + #[cfg(target_os = "linux")] + pub fn prepare_vulkan( + &mut self, + _ctx: &crate::context::vulkan::VulkanContext, + cmd: ash::vk::CommandBuffer, + slot: usize, + ) { + let Some(state) = self.vulkan.as_mut() else { + return; + }; + state.atlas_grayscale.flush_uploads( + &state.device, + &state.instance, + state.physical_device, + cmd, + slot, + ); + state.atlas_color.flush_uploads( + &state.device, + &state.instance, + state.physical_device, + cmd, + slot, + ); + } + + /// Record the UI text pass into `cmd`. Caller has already opened + /// the dynamic-rendering pass and set viewport/scissor. No-op + /// when no instances were recorded this frame or the Vulkan + /// state isn't initialised. + #[cfg(target_os = "linux")] + pub fn render_vulkan( + &mut self, + cmd: ash::vk::CommandBuffer, + slot: usize, + viewport: [f32; 2], + ) { + let instance_count = self.instances.len(); + if instance_count == 0 { + return; + } + let Some(state) = self.vulkan.as_mut() else { + return; + }; + + // Upload uniforms (viewport + 8B pad — std140). + let uniforms: [f32; 4] = [viewport[0], viewport[1], 0.0, 0.0]; + unsafe { + let dst = state.uniform_buffers[slot].as_mut_ptr() as *mut [f32; 4]; + std::ptr::write(dst, uniforms); + } + + // Grow per-slot instance buffer if needed. + let needed_bytes = instance_count * std::mem::size_of::(); + if instance_count > state.instance_capacity[slot] { + let new_cap = instance_count.next_power_of_two().max(256); + state.instance_buffers[slot] = + Some(crate::context::vulkan::allocate_host_visible_buffer_raw( + &state.device, + &state.instance, + state.physical_device, + (new_cap * std::mem::size_of::()) as u64, + ash::vk::BufferUsageFlags::VERTEX_BUFFER, + )); + state.instance_capacity[slot] = new_cap; + } + let instance_buf = state.instance_buffers[slot].as_ref().unwrap(); + unsafe { + std::ptr::copy_nonoverlapping( + self.instances.as_ptr() as *const u8, + instance_buf.as_mut_ptr(), + needed_bytes, + ); + } + + unsafe { + state.device.cmd_bind_pipeline( + cmd, + ash::vk::PipelineBindPoint::GRAPHICS, + state.pipeline, + ); + state.device.cmd_bind_descriptor_sets( + cmd, + ash::vk::PipelineBindPoint::GRAPHICS, + state.pipeline_layout, + 0, + &[ + state.uniform_descriptor_sets[slot], + state.atlas_descriptor_set, + ], + &[], + ); + state + .device + .cmd_bind_vertex_buffers(cmd, 0, &[instance_buf.handle()], &[0]); + state.device.cmd_draw(cmd, 4, instance_count as u32, 0, 0); + } + } } // Helpers @@ -1137,3 +1339,368 @@ fn premul_blend_wgpu() -> wgpu::BlendState { }, } } + +// Vulkan pipeline + state construction + +// Compiled at build time by `sugarloaf/build.rs`. The fragment +// shader is shared with the grid text pass — same atlas sampling, +// same inputs. +#[cfg(target_os = "linux")] +const UI_TEXT_VERT_SPV: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/ui_text.vert.spv")); +#[cfg(target_os = "linux")] +const UI_TEXT_FRAG_SPV: &[u8] = + include_bytes!(concat!(env!("OUT_DIR"), "/grid_text.frag.spv")); + +#[cfg(target_os = "linux")] +fn build_text_vulkan_state( + ctx: &crate::context::vulkan::VulkanContext, +) -> TextVulkanState { + use crate::context::vulkan::FRAMES_IN_FLIGHT; + use ash::vk; + + let device = ctx.device().clone(); + let instance = ctx.instance().clone(); + let physical_device = ctx.physical_device(); + + let atlas_grayscale = crate::grid::vulkan::VulkanGlyphAtlas::new_grayscale(ctx); + let atlas_color = crate::grid::vulkan::VulkanGlyphAtlas::new_color(ctx); + let sampler = create_text_sampler(&device); + + let uniform_buffers = std::array::from_fn(|_| { + ctx.allocate_host_visible_buffer(16, vk::BufferUsageFlags::UNIFORM_BUFFER) + }); + + // Layouts: set 0 = uniform (per slot), set 1 = atlases (shared). + let uniform_descriptor_set_layout = unsafe { + let bindings = [vk::DescriptorSetLayoutBinding::default() + .binding(0) + .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::VERTEX)]; + let info = vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings); + device + .create_descriptor_set_layout(&info, None) + .expect("create_descriptor_set_layout(ui_text uniform)") + }; + let atlas_descriptor_set_layout = unsafe { + let bindings = [ + vk::DescriptorSetLayoutBinding::default() + .binding(0) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::FRAGMENT), + vk::DescriptorSetLayoutBinding::default() + .binding(1) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .descriptor_count(1) + .stage_flags(vk::ShaderStageFlags::FRAGMENT), + ]; + let info = vk::DescriptorSetLayoutCreateInfo::default().bindings(&bindings); + device + .create_descriptor_set_layout(&info, None) + .expect("create_descriptor_set_layout(ui_text atlas)") + }; + + // Pool sized for FRAMES_IN_FLIGHT uniform sets + 1 atlas set. + let descriptor_pool = unsafe { + let sizes = [ + vk::DescriptorPoolSize { + ty: vk::DescriptorType::UNIFORM_BUFFER, + descriptor_count: FRAMES_IN_FLIGHT as u32, + }, + vk::DescriptorPoolSize { + ty: vk::DescriptorType::COMBINED_IMAGE_SAMPLER, + descriptor_count: 2, + }, + ]; + let info = vk::DescriptorPoolCreateInfo::default() + .max_sets((FRAMES_IN_FLIGHT + 1) as u32) + .pool_sizes(&sizes); + device + .create_descriptor_pool(&info, None) + .expect("create_descriptor_pool(ui_text)") + }; + + let uniform_descriptor_sets = unsafe { + let layouts = [uniform_descriptor_set_layout; FRAMES_IN_FLIGHT]; + let info = vk::DescriptorSetAllocateInfo::default() + .descriptor_pool(descriptor_pool) + .set_layouts(&layouts); + let sets = device + .allocate_descriptor_sets(&info) + .expect("allocate_descriptor_sets(ui_text uniform)"); + let mut out = [vk::DescriptorSet::null(); FRAMES_IN_FLIGHT]; + out.copy_from_slice(&sets); + out + }; + let atlas_descriptor_set = unsafe { + let layouts = [atlas_descriptor_set_layout]; + let info = vk::DescriptorSetAllocateInfo::default() + .descriptor_pool(descriptor_pool) + .set_layouts(&layouts); + device + .allocate_descriptor_sets(&info) + .expect("allocate_descriptor_sets(ui_text atlas)")[0] + }; + + // Update descriptor sets. + for slot in 0..FRAMES_IN_FLIGHT { + let uniform_info = vk::DescriptorBufferInfo::default() + .buffer(uniform_buffers[slot].handle()) + .offset(0) + .range(uniform_buffers[slot].size()); + let infos = [uniform_info]; + let write = vk::WriteDescriptorSet::default() + .dst_set(uniform_descriptor_sets[slot]) + .dst_binding(0) + .descriptor_type(vk::DescriptorType::UNIFORM_BUFFER) + .buffer_info(&infos); + unsafe { + device.update_descriptor_sets(&[write], &[]); + } + } + { + let gray_info = vk::DescriptorImageInfo::default() + .sampler(sampler) + .image_view(atlas_grayscale.image_view()) + .image_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL); + let gray_infos = [gray_info]; + let color_info = vk::DescriptorImageInfo::default() + .sampler(sampler) + .image_view(atlas_color.image_view()) + .image_layout(vk::ImageLayout::SHADER_READ_ONLY_OPTIMAL); + let color_infos = [color_info]; + let writes = [ + vk::WriteDescriptorSet::default() + .dst_set(atlas_descriptor_set) + .dst_binding(0) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .image_info(&gray_infos), + vk::WriteDescriptorSet::default() + .dst_set(atlas_descriptor_set) + .dst_binding(1) + .descriptor_type(vk::DescriptorType::COMBINED_IMAGE_SAMPLER) + .image_info(&color_infos), + ]; + unsafe { + device.update_descriptor_sets(&writes, &[]); + } + } + + let pipeline_layout = unsafe { + let set_layouts = [uniform_descriptor_set_layout, atlas_descriptor_set_layout]; + let info = vk::PipelineLayoutCreateInfo::default().set_layouts(&set_layouts); + device + .create_pipeline_layout(&info, None) + .expect("create_pipeline_layout(ui_text)") + }; + let pipeline = build_ui_text_pipeline_vulkan( + &device, + ctx.pipeline_cache(), + pipeline_layout, + ctx.swapchain_format(), + ); + + TextVulkanState { + device, + instance, + physical_device, + atlas_grayscale, + atlas_color, + sampler, + uniform_buffers, + instance_buffers: std::array::from_fn(|_| None), + instance_capacity: [0; FRAMES_IN_FLIGHT], + descriptor_pool, + uniform_descriptor_set_layout, + atlas_descriptor_set_layout, + uniform_descriptor_sets, + atlas_descriptor_set, + pipeline_layout, + pipeline, + } +} + +#[cfg(target_os = "linux")] +fn create_text_sampler(device: &ash::Device) -> ash::vk::Sampler { + use ash::vk; + let info = vk::SamplerCreateInfo::default() + .mag_filter(vk::Filter::NEAREST) + .min_filter(vk::Filter::NEAREST) + .mipmap_mode(vk::SamplerMipmapMode::NEAREST) + .address_mode_u(vk::SamplerAddressMode::CLAMP_TO_EDGE) + .address_mode_v(vk::SamplerAddressMode::CLAMP_TO_EDGE) + .address_mode_w(vk::SamplerAddressMode::CLAMP_TO_EDGE); + unsafe { + device + .create_sampler(&info, None) + .expect("create_sampler(ui_text)") + } +} + +#[cfg(target_os = "linux")] +fn build_ui_text_pipeline_vulkan( + device: &ash::Device, + pipeline_cache: ash::vk::PipelineCache, + layout: ash::vk::PipelineLayout, + color_format: ash::vk::Format, +) -> ash::vk::Pipeline { + use ash::vk; + + let vert = load_shader_module_vulkan(device, UI_TEXT_VERT_SPV); + let frag = load_shader_module_vulkan(device, UI_TEXT_FRAG_SPV); + let entry = c"main"; + let stages = [ + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::VERTEX) + .module(vert) + .name(entry), + vk::PipelineShaderStageCreateInfo::default() + .stage(vk::ShaderStageFlags::FRAGMENT) + .module(frag) + .name(entry), + ]; + + // Vertex input mirrors `TextInstance` (36 bytes). + let bindings = [vk::VertexInputBindingDescription::default() + .binding(0) + .stride(std::mem::size_of::() as u32) + .input_rate(vk::VertexInputRate::INSTANCE)]; + let attrs = [ + // 0: pos vec2 @ 0 + vk::VertexInputAttributeDescription::default() + .location(0) + .binding(0) + .format(vk::Format::R32G32_SFLOAT) + .offset(0), + // 1: glyph_pos uvec2 @ 8 + vk::VertexInputAttributeDescription::default() + .location(1) + .binding(0) + .format(vk::Format::R32G32_UINT) + .offset(8), + // 2: glyph_size uvec2 @ 16 + vk::VertexInputAttributeDescription::default() + .location(2) + .binding(0) + .format(vk::Format::R32G32_UINT) + .offset(16), + // 3: bearings ivec2 @ 24 + vk::VertexInputAttributeDescription::default() + .location(3) + .binding(0) + .format(vk::Format::R16G16_SINT) + .offset(24), + // 4: color vec4 @ 28 + vk::VertexInputAttributeDescription::default() + .location(4) + .binding(0) + .format(vk::Format::R8G8B8A8_UNORM) + .offset(28), + // 5: atlas u8 @ 32 + vk::VertexInputAttributeDescription::default() + .location(5) + .binding(0) + .format(vk::Format::R8_UINT) + .offset(32), + ]; + let vertex_input = vk::PipelineVertexInputStateCreateInfo::default() + .vertex_binding_descriptions(&bindings) + .vertex_attribute_descriptions(&attrs); + + let input_assembly = vk::PipelineInputAssemblyStateCreateInfo::default() + .topology(vk::PrimitiveTopology::TRIANGLE_STRIP); + let viewport_state = vk::PipelineViewportStateCreateInfo::default() + .viewport_count(1) + .scissor_count(1); + let rasterization = vk::PipelineRasterizationStateCreateInfo::default() + .polygon_mode(vk::PolygonMode::FILL) + .cull_mode(vk::CullModeFlags::NONE) + .front_face(vk::FrontFace::COUNTER_CLOCKWISE) + .line_width(1.0); + let multisample = vk::PipelineMultisampleStateCreateInfo::default() + .rasterization_samples(vk::SampleCountFlags::TYPE_1); + + // Premultiplied-over-from-one — fragment returns premultiplied. + let blend_attachment = vk::PipelineColorBlendAttachmentState::default() + .blend_enable(true) + .src_color_blend_factor(vk::BlendFactor::ONE) + .dst_color_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .color_blend_op(vk::BlendOp::ADD) + .src_alpha_blend_factor(vk::BlendFactor::ONE) + .dst_alpha_blend_factor(vk::BlendFactor::ONE_MINUS_SRC_ALPHA) + .alpha_blend_op(vk::BlendOp::ADD) + .color_write_mask(vk::ColorComponentFlags::RGBA); + let blend_attachments = [blend_attachment]; + let color_blend = + vk::PipelineColorBlendStateCreateInfo::default().attachments(&blend_attachments); + + let dynamic_states = [vk::DynamicState::VIEWPORT, vk::DynamicState::SCISSOR]; + let dynamic_state = + vk::PipelineDynamicStateCreateInfo::default().dynamic_states(&dynamic_states); + + let color_attachment_formats = [color_format]; + let mut rendering = vk::PipelineRenderingCreateInfo::default() + .color_attachment_formats(&color_attachment_formats); + + let pipeline_info = vk::GraphicsPipelineCreateInfo::default() + .stages(&stages) + .vertex_input_state(&vertex_input) + .input_assembly_state(&input_assembly) + .viewport_state(&viewport_state) + .rasterization_state(&rasterization) + .multisample_state(&multisample) + .color_blend_state(&color_blend) + .dynamic_state(&dynamic_state) + .layout(layout) + .push_next(&mut rendering); + + let pipeline = unsafe { + device + .create_graphics_pipelines(pipeline_cache, &[pipeline_info], None) + .map_err(|(_, e)| e) + .expect("create_graphics_pipelines(ui_text)")[0] + }; + unsafe { + device.destroy_shader_module(vert, None); + device.destroy_shader_module(frag, None); + } + pipeline +} + +#[cfg(target_os = "linux")] +fn load_shader_module_vulkan( + device: &ash::Device, + bytes: &[u8], +) -> ash::vk::ShaderModule { + use ash::vk; + let code = ash::util::read_spv(&mut std::io::Cursor::new(bytes)) + .expect("read_spv (embedded ui_text shader is valid)"); + let info = vk::ShaderModuleCreateInfo::default().code(&code); + unsafe { + device + .create_shader_module(&info, None) + .expect("create_shader_module(ui_text)") + } +} + +#[cfg(target_os = "linux")] +impl Drop for TextVulkanState { + fn drop(&mut self) { + unsafe { + let _ = self.device.device_wait_idle(); + self.device.destroy_pipeline(self.pipeline, None); + self.device + .destroy_pipeline_layout(self.pipeline_layout, None); + self.device + .destroy_descriptor_pool(self.descriptor_pool, None); + self.device + .destroy_descriptor_set_layout(self.atlas_descriptor_set_layout, None); + self.device + .destroy_descriptor_set_layout(self.uniform_descriptor_set_layout, None); + self.device.destroy_sampler(self.sampler, None); + // Buffers + atlas images drop themselves. + } + } +}