diff --git a/Cargo.toml b/Cargo.toml index 397ece2..5873f00 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -2,6 +2,8 @@ resolver = "2" members = [ + "crates/shaders", + "integrations/vello_svg", "examples/headless", diff --git a/crates/shaders/Cargo.toml b/crates/shaders/Cargo.toml new file mode 100644 index 0000000..d3969ef --- /dev/null +++ b/crates/shaders/Cargo.toml @@ -0,0 +1,19 @@ +[package] +name = "vello_shaders" +version = "0.1.0" +edition = "2021" + +[features] +default = ["compile", "wgsl", "msl"] +compile = ["naga", "thiserror"] +wgsl = [] +msl = [] + +[dependencies] +naga = { git = "https://github.com/gfx-rs/naga", rev = "53d62b9", features = ["wgsl-in", "msl-out", "validate"], optional = true } +thiserror = { version = "1.0.40", optional = true } + +[build-dependencies] +naga = { git = "https://github.com/gfx-rs/naga", rev = "53d62b9", features = ["wgsl-in", "msl-out", "validate"] } +thiserror = "1.0.40" + diff --git a/crates/shaders/README.md b/crates/shaders/README.md new file mode 100644 index 0000000..9bec768 --- /dev/null +++ b/crates/shaders/README.md @@ -0,0 +1,7 @@ +The `vello_shaders` crate provides a utility library to integrate the Vello shader modules into any +renderer project. The crate provides the necessary metadata to construct the individual compute +pipelines on any GPU API while leaving the responsibility of all API interactions (such as +resource management and command encoding) up to the client. + +The shaders can be pre-compiled to any target shading language at build time based on feature flags. +Currently only WGSL and Metal Shading Language are supported. diff --git a/crates/shaders/build.rs b/crates/shaders/build.rs new file mode 100644 index 0000000..7e5a952 --- /dev/null +++ b/crates/shaders/build.rs @@ -0,0 +1,107 @@ +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +#[path = "src/compile/mod.rs"] +mod compile; +#[path = "src/types.rs"] +mod types; + +use std::env; +use std::fmt::Write; +use std::path::Path; + +use compile::ShaderInfo; + +fn main() { + let out_dir = env::var_os("OUT_DIR").unwrap(); + let dest_path = Path::new(&out_dir).join("shaders.rs"); + let mut shaders = compile::ShaderInfo::from_dir("../../shader"); + // Drop the HashMap and sort by name so that we get deterministic order. + let mut shaders = shaders.drain().collect::>(); + shaders.sort_by(|x, y| x.0.cmp(&y.0)); + let mut buf = String::default(); + write_types(&mut buf, &shaders).unwrap(); + if cfg!(feature = "wgsl") { + write_shaders(&mut buf, "wgsl", &shaders, |info| { + info.source.as_bytes().to_owned() + }) + .unwrap(); + } + if cfg!(feature = "msl") { + write_shaders(&mut buf, "msl", &shaders, |info| { + compile::msl::translate(info).unwrap().as_bytes().to_owned() + }) + .unwrap(); + } + std::fs::write(&dest_path, &buf).unwrap(); + println!("cargo:rerun-if-changed=../shader"); +} + +fn write_types(buf: &mut String, shaders: &[(String, ShaderInfo)]) -> Result<(), std::fmt::Error> { + writeln!(buf, "pub struct Shaders<'a> {{")?; + for (name, _) in shaders { + writeln!(buf, " pub {name}: ComputeShader<'a>,")?; + } + writeln!(buf, "}}")?; + writeln!(buf, "pub struct Pipelines {{")?; + for (name, _) in shaders { + writeln!(buf, " pub {name}: T,")?; + } + writeln!(buf, "}}")?; + writeln!(buf, "impl Pipelines {{")?; + writeln!(buf, " pub fn from_shaders>(shaders: &Shaders, device: &H::Device, host: &mut H) -> Result {{")?; + writeln!(buf, " Ok(Self {{")?; + for (name, _) in shaders { + writeln!( + buf, + " {name}: host.new_compute_pipeline(device, &shaders.{name})?," + )?; + } + writeln!(buf, " }})")?; + writeln!(buf, " }}")?; + writeln!(buf, "}}")?; + Ok(()) +} + +fn write_shaders( + buf: &mut String, + mod_name: &str, + shaders: &[(String, ShaderInfo)], + translate: impl Fn(&ShaderInfo) -> Vec, +) -> Result<(), std::fmt::Error> { + writeln!(buf, "pub mod {mod_name} {{")?; + writeln!(buf, " use super::*;")?; + writeln!(buf, " use BindType::*;")?; + writeln!(buf, " pub const SHADERS: Shaders<'static> = Shaders {{")?; + for (name, info) in shaders { + let bind_tys = info + .bindings + .iter() + .map(|binding| binding.ty) + .collect::>(); + let wg_bufs = &info.workgroup_buffers; + let source = translate(info); + writeln!(buf, " {name}: ComputeShader {{")?; + writeln!(buf, " name: Cow::Borrowed({:?}),", name)?; + writeln!( + buf, + " code: Cow::Borrowed(&{:?}),", + source.as_slice() + )?; + writeln!( + buf, + " workgroup_size: {:?},", + info.workgroup_size + )?; + writeln!(buf, " bindings: Cow::Borrowed(&{:?}),", bind_tys)?; + writeln!( + buf, + " workgroup_buffers: Cow::Borrowed(&{:?}),", + wg_bufs + )?; + writeln!(buf, " }},")?; + } + writeln!(buf, " }};")?; + writeln!(buf, "}}")?; + Ok(()) +} diff --git a/crates/shaders/src/compile/mod.rs b/crates/shaders/src/compile/mod.rs new file mode 100644 index 0000000..74cafda --- /dev/null +++ b/crates/shaders/src/compile/mod.rs @@ -0,0 +1,197 @@ +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +use { + naga::{ + front::wgsl, + valid::{Capabilities, ModuleInfo, ValidationError, ValidationFlags}, + AddressSpace, ArraySize, ConstantInner, ImageClass, Module, ScalarValue, StorageAccess, + WithSpan, + }, + std::{ + collections::{HashMap, HashSet}, + path::Path, + }, + thiserror::Error, +}; + +pub mod permutations; +pub mod preprocess; + +pub mod msl; + +use crate::types::{BindType, BindingInfo, WorkgroupBufferInfo}; + +#[derive(Error, Debug)] +pub enum Error { + #[error("failed to parse shader: {0}")] + Parse(#[from] wgsl::ParseError), + + #[error("failed to validate shader: {0}")] + Validate(#[from] WithSpan), + + #[error("missing entry point function")] + EntryPointNotFound, +} + +#[derive(Debug)] +pub struct ShaderInfo { + pub source: String, + pub module: Module, + pub module_info: ModuleInfo, + pub workgroup_size: [u32; 3], + pub bindings: Vec, + pub workgroup_buffers: Vec, +} + +impl ShaderInfo { + pub fn new(source: String, entry_point: &str) -> Result { + let module = wgsl::parse_str(&source)?; + let module_info = naga::valid::Validator::new( + ValidationFlags::all() & !ValidationFlags::CONTROL_FLOW_UNIFORMITY, + Capabilities::all(), + ) + .validate(&module)?; + let (entry_index, entry) = module + .entry_points + .iter() + .enumerate() + .find(|(_, entry)| entry.name.as_str() == entry_point) + .ok_or(Error::EntryPointNotFound)?; + let mut bindings = vec![]; + let mut workgroup_buffers = vec![]; + let mut wg_buffer_idx = 0; + let entry_info = module_info.get_entry_point(entry_index); + for (var_handle, var) in module.global_variables.iter() { + if entry_info[var_handle].is_empty() { + continue; + } + let binding_ty = match module.types[var.ty].inner { + naga::TypeInner::BindingArray { base, .. } => &module.types[base].inner, + ref ty => ty, + }; + let Some(binding) = &var.binding else { + if var.space == AddressSpace::WorkGroup { + let index = wg_buffer_idx; + wg_buffer_idx += 1; + let size_in_bytes = match binding_ty { + naga::TypeInner::Array { + size: ArraySize::Constant(const_handle), + stride, + .. + } => { + let size: u32 = match module.constants[*const_handle].inner { + ConstantInner::Scalar { value, width: _ } => match value { + ScalarValue::Uint(value) => value.try_into().unwrap(), + ScalarValue::Sint(value) => value.try_into().unwrap(), + _ => continue, + }, + ConstantInner::Composite { .. } => continue, + }; + size * stride + }, + naga::TypeInner::Struct { span, .. } => *span, + naga::TypeInner::Scalar { width, ..} => *width as u32, + naga::TypeInner::Vector { width, ..} => *width as u32, + naga::TypeInner::Matrix { width, ..} => *width as u32, + naga::TypeInner::Atomic { width, ..} => *width as u32, + _ => { + // Not a valid workgroup variable type. At least not one that is used + // in our shaders. + continue; + } + }; + workgroup_buffers.push(WorkgroupBufferInfo { + size_in_bytes, + index, + }); + } + continue; + }; + let mut resource = BindingInfo { + name: var.name.clone(), + location: (binding.group, binding.binding), + ty: BindType::Buffer, + }; + if let naga::TypeInner::Image { class, .. } = &binding_ty { + resource.ty = BindType::ImageRead; + if let ImageClass::Storage { access, .. } = class { + if access.contains(StorageAccess::STORE) { + resource.ty = BindType::Image; + } + } + } else { + resource.ty = BindType::BufReadOnly; + match var.space { + AddressSpace::Storage { access } => { + if access.contains(StorageAccess::STORE) { + resource.ty = BindType::Buffer; + } + } + AddressSpace::Uniform => { + resource.ty = BindType::Uniform; + } + _ => {} + } + } + bindings.push(resource); + } + bindings.sort_by_key(|res| res.location); + let workgroup_size = entry.workgroup_size; + Ok(ShaderInfo { + source, + module, + module_info, + workgroup_size, + bindings, + workgroup_buffers, + }) + } + + pub fn from_dir(shader_dir: impl AsRef) -> HashMap { + use std::fs; + let shader_dir = shader_dir.as_ref(); + let permutation_map = if let Ok(permutations_source) = + std::fs::read_to_string(shader_dir.join("permutations")) + { + permutations::parse(&permutations_source) + } else { + Default::default() + }; + println!("{:?}", permutation_map); + let imports = preprocess::get_imports(shader_dir); + let mut info = HashMap::default(); + let mut defines = HashSet::default(); + defines.insert("full".to_string()); + for entry in shader_dir + .read_dir() + .expect("Can read shader import directory") + { + let entry = entry.expect("Can continue reading shader import directory"); + if entry.file_type().unwrap().is_file() { + let file_name = entry.file_name(); + if let Some(name) = file_name.to_str() { + let suffix = ".wgsl"; + if let Some(shader_name) = name.strip_suffix(suffix) { + let contents = fs::read_to_string(shader_dir.join(&file_name)) + .expect("Could read shader {shader_name} contents"); + if let Some(permutations) = permutation_map.get(shader_name) { + for permutation in permutations { + let mut defines = defines.clone(); + defines.extend(permutation.defines.iter().cloned()); + let source = preprocess::preprocess(&contents, &defines, &imports); + let shader_info = Self::new(source.clone(), "main").unwrap(); + info.insert(permutation.name.clone(), shader_info); + } + } else { + let source = preprocess::preprocess(&contents, &defines, &imports); + let shader_info = Self::new(source.clone(), "main").unwrap(); + info.insert(shader_name.to_string(), shader_info); + } + } + } + } + } + info + } +} diff --git a/crates/shaders/src/compile/msl.rs b/crates/shaders/src/compile/msl.rs new file mode 100644 index 0000000..5f7b831 --- /dev/null +++ b/crates/shaders/src/compile/msl.rs @@ -0,0 +1,56 @@ +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +use naga::back::msl; + +use super::{BindType, ShaderInfo}; + +pub fn translate(shader: &ShaderInfo) -> Result { + let mut map = msl::EntryPointResourceMap::default(); + let mut buffer_index = 0u8; + let mut image_index = 0u8; + let mut binding_map = msl::BindingMap::default(); + for resource in &shader.bindings { + let binding = naga::ResourceBinding { + group: resource.location.0, + binding: resource.location.1, + }; + let mut target = msl::BindTarget::default(); + match resource.ty { + BindType::Buffer | BindType::BufReadOnly | BindType::Uniform => { + target.buffer = Some(buffer_index); + buffer_index += 1; + } + BindType::Image | BindType::ImageRead => { + target.texture = Some(image_index); + image_index += 1; + } + } + target.mutable = resource.ty.is_mutable(); + binding_map.insert(binding, target); + } + map.insert( + "main".to_string(), + msl::EntryPointResources { + resources: binding_map, + push_constant_buffer: None, + sizes_buffer: Some(30), + }, + ); + let options = msl::Options { + lang_version: (2, 0), + per_entry_point_map: map, + inline_samplers: vec![], + spirv_cross_compatibility: false, + fake_missing_bindings: false, + bounds_check_policies: naga::proc::BoundsCheckPolicies::default(), + zero_initialize_workgroup_memory: false, + }; + let (source, _) = msl::write_string( + &shader.module, + &shader.module_info, + &options, + &msl::PipelineOptions::default(), + )?; + Ok(source) +} diff --git a/crates/shaders/src/compile/permutations.rs b/crates/shaders/src/compile/permutations.rs new file mode 100644 index 0000000..f85f667 --- /dev/null +++ b/crates/shaders/src/compile/permutations.rs @@ -0,0 +1,44 @@ +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +use std::collections::HashMap; + +#[derive(Debug)] +pub struct Permutation { + /// The new name for the permutation + pub name: String, + /// Set of defines to apply for the permutation + pub defines: Vec, +} + +pub fn parse(source: &str) -> HashMap> { + let mut map: HashMap> = Default::default(); + let mut current_source: Option = None; + for line in source.lines() { + let line = line.trim(); + if line.is_empty() || line.starts_with('#') { + continue; + } + if let Some(line) = line.strip_prefix('+') { + if let Some(source) = ¤t_source { + let mut parts = line.split(':').map(|s| s.trim()); + let Some(name) = parts.next() else { + continue; + }; + let mut defines = vec![]; + if let Some(define_list) = parts.next() { + defines.extend(define_list.split(' ').map(|s| s.trim().to_string())); + } + map.entry(source.to_string()) + .or_default() + .push(Permutation { + name: name.to_string(), + defines, + }); + } + } else { + current_source = Some(line.to_string()); + } + } + map +} diff --git a/crates/shaders/src/compile/preprocess.rs b/crates/shaders/src/compile/preprocess.rs new file mode 100644 index 0000000..917f83f --- /dev/null +++ b/crates/shaders/src/compile/preprocess.rs @@ -0,0 +1,162 @@ +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +use std::{ + collections::{HashMap, HashSet}, + fs, + path::Path, + vec, +}; + +pub fn get_imports(shader_dir: &Path) -> HashMap { + let mut imports = HashMap::new(); + let imports_dir = shader_dir.join("shared"); + for entry in imports_dir + .read_dir() + .expect("Can read shader import directory") + { + let entry = entry.expect("Can continue reading shader import directory"); + if entry.file_type().unwrap().is_file() { + let file_name = entry.file_name(); + if let Some(name) = file_name.to_str() { + let suffix = ".wgsl"; + if let Some(import_name) = name.strip_suffix(suffix) { + let contents = fs::read_to_string(imports_dir.join(&file_name)) + .expect("Could read shader {import_name} contents"); + imports.insert(import_name.to_owned(), contents); + } + } + } + } + imports +} + +pub struct StackItem { + active: bool, + else_passed: bool, +} + +pub fn preprocess( + input: &str, + defines: &HashSet, + imports: &HashMap, +) -> String { + let mut output = String::with_capacity(input.len()); + let mut stack = vec![]; + 'all_lines: for (line_number, mut line) in input.lines().enumerate() { + loop { + if line.is_empty() { + break; + } + let hash_index = line.find('#'); + let comment_index = line.find("//"); + let hash_index = match (hash_index, comment_index) { + (Some(hash_index), None) => hash_index, + (Some(hash_index), Some(comment_index)) if hash_index < comment_index => hash_index, + // Add this line to the output - all directives are commented out or there are no directives + _ => break, + }; + let directive_start = &line[hash_index + '#'.len_utf8()..]; + let directive_len = directive_start + // The first character which can't be part of the directive name marks the end of the directive + // In practise this should always be whitespace, but in theory a 'unit' directive + // could be added + .find(|c: char| !c.is_alphanumeric()) + .unwrap_or(directive_start.len()); + let directive = &directive_start[..directive_len]; + let directive_is_at_start = line.trim_start().starts_with('#'); + + match directive { + if_item @ ("ifdef" | "ifndef" | "else" | "endif") if !directive_is_at_start => { + eprintln!("#{if_item} directives must be the first non_whitespace items on their line, ignoring (line {line_number})"); + break; + } + def_test @ ("ifdef" | "ifndef") => { + let def = directive_start[directive_len..].trim(); + let exists = defines.contains(def); + let mode = def_test == "ifdef"; + stack.push(StackItem { + active: mode == exists, + else_passed: false, + }); + // Don't add this line to the output; instead process the next line + continue 'all_lines; + } + "else" => { + let item = stack.last_mut(); + if let Some(item) = item { + if item.else_passed { + eprintln!("Second else for same ifdef/ifndef (line {line_number}); ignoring second else") + } else { + item.else_passed = true; + item.active = !item.active; + } + } + let remainder = directive_start[directive_len..].trim(); + if !remainder.is_empty() { + eprintln!("#else directives don't take an argument. `{remainder}` will not be in output (line {line_number})"); + } + // Don't add this line to the output; it should be empty (see warning above) + continue 'all_lines; + } + "endif" => { + if stack.pop().is_none() { + eprintln!("Mismatched endif (line {line_number})"); + } + let remainder = directive_start[directive_len..].trim(); + if !remainder.is_empty() { + eprintln!("#endif directives don't take an argument. `{remainder}` will not be in output (line {line_number})"); + } + // Don't add this line to the output; it should be empty (see warning above) + continue 'all_lines; + } + "import" => { + output.push_str(&line[..hash_index]); + let directive_end = &directive_start[directive_len..]; + let import_name_start = if let Some(import_name_start) = + directive_end.find(|c: char| !c.is_whitespace()) + { + import_name_start + } else { + eprintln!("#import needs a non_whitespace argument (line {line_number})"); + continue 'all_lines; + }; + let import_name_start = &directive_end[import_name_start..]; + let import_name_end_index = import_name_start + // The first character which can't be part of the import name marks the end of the import + .find(|c: char| !(c == '_' || c.is_alphanumeric())) + .unwrap_or(import_name_start.len()); + let import_name = &import_name_start[..import_name_end_index]; + line = &import_name_start[import_name_end_index..]; + let import = imports.get(import_name); + if let Some(import) = import { + // In theory, we can cache this until the top item of the stack changes + // However, in practise there will only ever be at most 2 stack items, so it's reasonable to just recompute it every time + if stack.iter().all(|item| item.active) { + output.push_str(&preprocess(import, defines, imports)); + } + } else { + eprintln!("Unknown import `{import_name}` (line {line_number})"); + } + continue; + } + val => { + eprintln!("Unknown preprocessor directive `{val}` (line {line_number})"); + } + } + } + if stack.iter().all(|item| item.active) { + // Naga does not yet recognize `const` but web does not allow global `let`. We + // use `let` in our canonical sources to satisfy wgsl-analyzer but replace with + // `const` when targeting web. + if line.starts_with("let ") { + output.push_str("const"); + output.push_str(&line[3..]); + } else { + output.push_str(line); + } + output.push('\n'); + } + } + output +} diff --git a/crates/shaders/src/lib.rs b/crates/shaders/src/lib.rs new file mode 100644 index 0000000..66ac937 --- /dev/null +++ b/crates/shaders/src/lib.rs @@ -0,0 +1,34 @@ +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +mod types; + +#[cfg(feature = "compile")] +pub mod compile; + +pub use types::{BindType, BindingInfo, WorkgroupBufferInfo}; + +use std::borrow::Cow; + +#[derive(Clone, Debug)] +pub struct ComputeShader<'a> { + pub name: Cow<'a, str>, + pub code: Cow<'a, [u8]>, + pub workgroup_size: [u32; 3], + pub bindings: Cow<'a, [BindType]>, + pub workgroup_buffers: Cow<'a, [WorkgroupBufferInfo]>, +} + +pub trait PipelineHost { + type Device; + type ComputePipeline; + type Error; + + fn new_compute_pipeline( + &mut self, + device: &Self::Device, + shader: &ComputeShader, + ) -> Result; +} + +include!(concat!(env!("OUT_DIR"), "/shaders.rs")); diff --git a/crates/shaders/src/types.rs b/crates/shaders/src/types.rs new file mode 100644 index 0000000..548f7e9 --- /dev/null +++ b/crates/shaders/src/types.rs @@ -0,0 +1,40 @@ +// Copyright 2023 The Vello authors +// SPDX-License-Identifier: Apache-2.0 OR MIT + +//! Types that are shared between the main crate and build. + +/// The type of resource that will be bound to a slot in a shader. +#[derive(Copy, Clone, PartialEq, Eq, Debug)] +pub enum BindType { + /// A storage buffer with read/write access. + Buffer, + /// A storage buffer with read only access. + BufReadOnly, + /// A small storage buffer to be used as uniforms. + Uniform, + /// A storage image. + Image, + /// A storage image with read only access. + ImageRead, + // TODO: Sampler, maybe others +} + +impl BindType { + pub fn is_mutable(self) -> bool { + matches!(self, Self::Buffer | Self::Image) + } +} + +#[derive(Clone, Debug)] +pub struct BindingInfo { + pub name: Option, + pub location: (u32, u32), + pub ty: BindType, +} + +#[derive(Clone, Debug)] +pub struct WorkgroupBufferInfo { + pub size_in_bytes: u32, + /// The order in which the workgroup variable is declared in the shader module. + pub index: u32, +} diff --git a/shader/permutations b/shader/permutations new file mode 100644 index 0000000..beb1e48 --- /dev/null +++ b/shader/permutations @@ -0,0 +1,3 @@ +pathtag_scan ++ pathtag_scan_large ++ pathtag_scan_small: small