diff --git a/Cargo.toml b/Cargo.toml index ca412e4e92..90b3dedd1f 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -39,6 +39,7 @@ deserialize = ["serde"] spv-in = ["petgraph", "spirv"] spv-out = ["spirv"] wgsl-in = ["codespan-reporting"] +wgsl-out = [] hlsl-out = [] [[bin]] diff --git a/README.md b/README.md index 5873d8effa..3ca6c920f3 100644 --- a/README.md +++ b/README.md @@ -21,7 +21,7 @@ Rust | | | | Back-end | Status | Feature | Notes | --------------- | ------------------ | -------- | ----- | SPIR-V | :white_check_mark: | spv-out | | -WGSL | | | | +WGSL | :construction: | wgsl-out | | Metal | :white_check_mark: | msl-out | | HLSL | :construction: | hlsl-out | | GLSL | :ok: | glsl-out | | diff --git a/bin/convert.rs b/bin/convert.rs index 492fa783c0..151fc29674 100644 --- a/bin/convert.rs +++ b/bin/convert.rs @@ -268,6 +268,13 @@ fn main() { let hlsl = hlsl::write_string(&module).unwrap_pretty(); fs::write(output_path, hlsl).unwrap(); } + #[cfg(feature = "wgsl-out")] + "wgsl" => { + use naga::back::wgsl; + + let wgsl = wgsl::write_string(&module).unwrap_pretty(); + fs::write(output_path, wgsl).unwrap(); + } other => { let _ = params; panic!( diff --git a/src/back/mod.rs b/src/back/mod.rs index efffb40a45..a478d971d4 100644 --- a/src/back/mod.rs +++ b/src/back/mod.rs @@ -10,6 +10,8 @@ pub mod hlsl; pub mod msl; #[cfg(feature = "spv-out")] pub mod spv; +#[cfg(feature = "wgsl-out")] +pub mod wgsl; impl crate::Expression { /// Returns the ref count, upon reaching which this expression diff --git a/src/back/wgsl/keywords.rs b/src/back/wgsl/keywords.rs new file mode 100644 index 0000000000..a0ec86810f --- /dev/null +++ b/src/back/wgsl/keywords.rs @@ -0,0 +1,126 @@ +// https://gpuweb.github.io/gpuweb/wgsl/#keyword-summary +pub const RESERVED: &[&str] = &[ + // Type-defining keywords + "ARRAY", + "BOOL", + "FLOAT32", + "INT32", + "MAT2x2", + "MAT2x3", + "MAT2x4", + "MAT3x2", + "MAT3x3", + "MAT3x4", + "MAT4x2", + "MAT4x3", + "MAT4x4", + "POINTER", + "SAMPLER", + "SAMPLER_COMPARISON", + "STRUCT", + "TEXTURE_1D", + "TEXTURE_2D", + "TEXTURE_2D_ARRAY", + "TEXTURE_3D", + "TEXTURE_CUBE", + "TEXTURE_CUBE_ARRAY", + "TEXTURE_MULTISAMPLED_2D", + "TEXTURE_STORAGE_1D", + "TEXTURE_STORAGE_2D", + "TEXTURE_STORAGE_2D_ARRAY", + "TEXTURE_STORAGE_3D", + "TEXTURE_DEPTH_2D", + "TEXTURE_DEPTH_2D_ARRAY", + "TEXTURE_DEPTH_CUBE", + "TEXTURE_DEPTH_CUBE_ARRAY", + "UINT32", + "VEC2", + "VEC3", + "VEC4", + // Other keywords + "BITCAST", + "BLOCK", + "BREAK", + "CASE", + "CONTINUE", + "CONTINUING", + "DEFAULT", + "DISCARD", + "ELSE", + "ELSE_IF", + "ENABLE", + "FALLTHROUGH", + "FALSE", + "FN", + "FOR", + "FUNCTION", + "IF", + "LET", + "LOOP", + "PRIVATE", + "RETURN", + "STORAGE", + "SWITCH", + "TRUE", + "TYPE", + "UNIFORM", + "VAR", + "WORKGROUP", + // Image format keywords + "R8UNORM", + "R8SNORM", + "R8UINT", + "R8SINT", + "R16UINT", + "R16SINT", + "R16FLOAT", + "RG8UNORM", + "RG8SNORM", + "RG8UINT", + "RG8SINT", + "R32UINT", + "R32SINT", + "R32FLOAT", + "RG16UINT", + "RG16SINT", + "RG16FLOAT", + "RGBA8UNORM", + "RGBA8UNORM-SRGB", + "RGBA8SNORM", + "RGBA8UINT", + "RGBA8SINT", + "BGRA8UNORM", + "BGRA8UNORM-SRGB", + "RGB10A2UNORM", + "RG11B10FLOAT", + "RG32UINT", + "RG32SINT", + "RG32FLOAT", + "RGBA16UINT", + "RGBA16SINT", + "RGBA16FLOAT", + "RGBA32UINT", + "RGBA32SINT", + "RGBA32FLOAT", + // Reserved Keywords + "asm", + "bf16", + "do", + "enum", + "f16", + "f64", + "i8", + "i16", + "i64", + "const", + "typedef", + "u8", + "u16", + "u64", + "unless", + "using", + "while", + "regardless", + "premerge", + "handle", +]; diff --git a/src/back/wgsl/mod.rs b/src/back/wgsl/mod.rs new file mode 100644 index 0000000000..e785b57a12 --- /dev/null +++ b/src/back/wgsl/mod.rs @@ -0,0 +1,23 @@ +mod keywords; +mod writer; + +use std::io::Error as IoError; +use std::string::FromUtf8Error; +use thiserror::Error; + +pub use writer::Writer; + +#[derive(Error, Debug)] +pub enum Error { + #[error(transparent)] + IoError(#[from] IoError), + #[error(transparent)] + Utf8(#[from] FromUtf8Error), +} + +pub fn write_string(module: &crate::Module) -> Result { + let mut w = Writer::new(Vec::new()); + w.write(module)?; + let output = String::from_utf8(w.finish())?; + Ok(output) +} diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs new file mode 100644 index 0000000000..206e96d11e --- /dev/null +++ b/src/back/wgsl/writer.rs @@ -0,0 +1,144 @@ +use super::Error; +use crate::FastHashMap; +use crate::{ + back::wgsl::keywords::RESERVED, Function, Module, ShaderStage, StructLevel, TypeInner, +}; +use crate::{ + proc::{NameKey, Namer}, + StructMember, +}; +use std::io::Write; + +const _INDENT: &str = " "; + +/// Shorthand result used internally by the backend +type BackendResult = Result<(), Error>; + +enum Decoration { + VertexStage, + FragmentStage, + ComputeStage { workgroup_size: [u32; 3] }, + Block, +} + +pub struct Writer { + out: W, + names: FastHashMap, + namer: Namer, +} + +impl Writer { + pub fn new(out: W) -> Self { + Writer { + out, + names: FastHashMap::default(), + namer: Namer::default(), + } + } + + pub fn write(&mut self, module: &Module) -> BackendResult { + self.names.clear(); + self.namer.reset(module, RESERVED, &mut self.names); + + // Write all structs + for (handle, ty) in module.types.iter() { + if let TypeInner::Struct { + level, ref members, .. + } = ty.inner + { + let name = &self.names[&NameKey::Type(handle)].clone(); + let block = level == StructLevel::Root; + self.write_struct(name, block, members)?; + writeln!(self.out)?; + } + } + + for (_, ep) in module.entry_points.iter().enumerate() { + let decoration = match ep.stage { + ShaderStage::Vertex => Decoration::VertexStage, + ShaderStage::Fragment => Decoration::FragmentStage, + ShaderStage::Compute => Decoration::ComputeStage { + workgroup_size: ep.workgroup_size, + }, + }; + + self.write_decoration(decoration)?; + // Add a newline after decoration + writeln!(self.out)?; + self.write_function(&ep.function)?; + } + + // Add a newline at the end of file + writeln!(self.out)?; + + Ok(()) + } + + /// Helper method used to write structs + /// https://gpuweb.github.io/gpuweb/wgsl/#functions + /// + /// # Notes + /// Ends in a newline + fn write_function(&mut self, func: &Function) -> BackendResult { + write!(self.out, "fn {}(", func.name.as_ref().unwrap())?; // TODO: unnamed function? + write!(self.out, ")")?; + + write!(self.out, "{{")?; + write!(self.out, "}}")?; + Ok(()) + } + + /// Helper method to write a decoration + /// + /// # Notes + /// Adds no leading or trailing whitespace + fn write_decoration(&mut self, decoration: Decoration) -> BackendResult { + write!(self.out, "[[")?; + match decoration { + Decoration::VertexStage => write!(self.out, "stage(vertex)")?, + Decoration::FragmentStage => write!(self.out, "stage(fragment)")?, + Decoration::ComputeStage { workgroup_size } => { + write!( + self.out, + "{}", + format!( + "stage(compute), workgroup_size({}, {}, {})", + workgroup_size[0], workgroup_size[1], workgroup_size[2] + ) + )?; + } + Decoration::Block => { + write!(self.out, "block")?; + } + }; + write!(self.out, "]]")?; + + Ok(()) + } + + /// Helper method used to write structs + /// + /// # Notes + /// Ends in a newline + fn write_struct( + &mut self, + name: &str, + block: bool, + _members: &[StructMember], + ) -> BackendResult { + if block { + self.write_decoration(Decoration::Block)?; + writeln!(self.out)?; + } + write!(self.out, "struct {} {{", name)?; + write!(self.out, "}}")?; + + writeln!(self.out)?; + + Ok(()) + } + + pub fn finish(self) -> W { + self.out + } +}