diff options
Diffstat (limited to 'third_party/rust/naga/src/back')
| -rw-r--r-- | third_party/rust/naga/src/back/glsl.rs | 1579 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/glsl/keywords.rs | 204 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/mod.rs | 8 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/msl/keywords.rs | 102 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/msl/mod.rs | 211 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/msl/writer.rs | 990 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/spv/helpers.rs | 20 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/spv/instructions.rs | 708 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/spv/layout.rs | 91 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/spv/layout_tests.rs | 166 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/spv/mod.rs | 52 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/spv/test_framework.rs | 27 | ||||
| -rw-r--r-- | third_party/rust/naga/src/back/spv/writer.rs | 1776 |
13 files changed, 5934 insertions, 0 deletions
diff --git a/third_party/rust/naga/src/back/glsl.rs b/third_party/rust/naga/src/back/glsl.rs new file mode 100644 index 0000000000..6a497a675b --- /dev/null +++ b/third_party/rust/naga/src/back/glsl.rs @@ -0,0 +1,1579 @@ +//! OpenGL shading language backend +//! +//! The main structure is [`Writer`](struct.Writer.html), it maintains internal state that is used +//! to output a `Module` into glsl +//! +//! # Supported versions +//! ### Core +//! - 330 +//! - 400 +//! - 410 +//! - 420 +//! - 430 +//! - 450 +//! - 460 +//! +//! ### ES +//! - 300 +//! - 310 +//! + +use crate::{ + proc::{ + CallGraph, CallGraphBuilder, Interface, NameKey, Namer, ResolveContext, ResolveError, + Typifier, Visitor, + }, + Arena, ArraySize, BinaryOperator, BuiltIn, ConservativeDepth, Constant, ConstantInner, + DerivativeAxis, Expression, FastHashMap, Function, FunctionOrigin, GlobalVariable, Handle, + ImageClass, Interpolation, IntrinsicFunction, LocalVariable, Module, ScalarKind, ShaderStage, + Statement, StorageAccess, StorageClass, StorageFormat, StructMember, Type, TypeInner, + UnaryOperator, +}; +use std::{ + cmp::Ordering, + fmt::{self, Error as FmtError}, + io::{Error as IoError, Write}, +}; + +/// Contains a constant with a slice of all the reserved keywords RESERVED_KEYWORDS +mod keywords; + +const SUPPORTED_CORE_VERSIONS: &[u16] = &[330, 400, 410, 420, 430, 440, 450]; +const SUPPORTED_ES_VERSIONS: &[u16] = &[300, 310]; + +#[derive(Debug)] +pub enum Error { + FormatError(FmtError), + IoError(IoError), + Type(ResolveError), + Custom(String), +} + +impl From<FmtError> for Error { + fn from(err: FmtError) -> Self { + Error::FormatError(err) + } +} + +impl From<IoError> for Error { + fn from(err: IoError) -> Self { + Error::IoError(err) + } +} + +impl From<ResolveError> for Error { + fn from(err: ResolveError) -> Self { + Error::Type(err) + } +} + +impl fmt::Display for Error { + fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result { + match self { + Error::FormatError(err) => write!(f, "Formatting error {}", err), + Error::IoError(err) => write!(f, "Io error: {}", err), + Error::Type(err) => write!(f, "Type error: {:?}", err), + Error::Custom(err) => write!(f, "{}", err), + } + } +} + +#[derive(Debug, Copy, Clone, PartialEq)] +pub enum Version { + Desktop(u16), + Embedded(u16), +} + +impl Version { + fn is_es(&self) -> bool { + match self { + Version::Desktop(_) => false, + Version::Embedded(_) => true, + } + } + + fn is_supported(&self) -> bool { + match self { + Version::Desktop(v) => SUPPORTED_CORE_VERSIONS.contains(v), + Version::Embedded(v) => SUPPORTED_ES_VERSIONS.contains(v), + } + } +} + +impl PartialOrd for Version { + fn partial_cmp(&self, other: &Self) -> Option<Ordering> { + match (*self, *other) { + (Version::Desktop(x), Version::Desktop(y)) => Some(x.cmp(&y)), + (Version::Embedded(x), Version::Embedded(y)) => Some(x.cmp(&y)), + _ => None, + } + } +} + +impl fmt::Display for Version { + fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result { + match self { + Version::Desktop(v) => write!(f, "{} core", v), + Version::Embedded(v) => write!(f, "{} es", v), + } + } +} + +#[derive(Debug, Clone)] +pub struct Options { + pub version: Version, + pub entry_point: (ShaderStage, String), +} + +#[derive(Debug, Clone)] +pub struct TextureMapping { + pub texture: Handle<GlobalVariable>, + pub sampler: Option<Handle<GlobalVariable>>, +} + +bitflags::bitflags! { + struct Features: u32 { + const BUFFER_STORAGE = 1; + const ARRAY_OF_ARRAYS = 1 << 1; + const DOUBLE_TYPE = 1 << 2; + const FULL_IMAGE_FORMATS = 1 << 3; + const MULTISAMPLED_TEXTURES = 1 << 4; + const MULTISAMPLED_TEXTURE_ARRAYS = 1 << 5; + const CUBE_TEXTURES_ARRAY = 1 << 6; + const COMPUTE_SHADER = 1 << 7; + const IMAGE_LOAD_STORE = 1 << 8; + const CONSERVATIVE_DEPTH = 1 << 9; + const TEXTURE_1D = 1 << 10; + const PUSH_CONSTANT = 1 << 11; + } +} + +struct FeaturesManager(Features); + +impl FeaturesManager { + pub fn new() -> Self { + Self(Features::empty()) + } + + pub fn request(&mut self, features: Features) { + self.0 |= features + } + + #[allow(clippy::collapsible_if)] + pub fn write(&self, version: Version, mut out: impl Write) -> Result<(), Error> { + if self.0.contains(Features::COMPUTE_SHADER) { + if version < Version::Embedded(310) || version < Version::Desktop(420) { + return Err(Error::Custom(format!( + "Version {} doesn't support compute shaders", + version + ))); + } + + if !version.is_es() { + // https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_compute_shader.txt + writeln!(out, "#extension GL_ARB_compute_shader : require")?; + } + } + + if self.0.contains(Features::BUFFER_STORAGE) { + if version < Version::Embedded(310) || version < Version::Desktop(400) { + return Err(Error::Custom(format!( + "Version {} doesn't support buffer storage class", + version + ))); + } + + if let Version::Desktop(_) = version { + // https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_shader_storage_buffer_object.txt + writeln!( + out, + "#extension GL_ARB_shader_storage_buffer_object : require" + )?; + } + } + + if self.0.contains(Features::DOUBLE_TYPE) { + if version.is_es() || version < Version::Desktop(150) { + return Err(Error::Custom(format!( + "Version {} doesn't support doubles", + version + ))); + } + + if version < Version::Desktop(400) { + // https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_gpu_shader_fp64.txt + writeln!(out, "#extension GL_ARB_gpu_shader_fp64 : require")?; + } + } + + if self.0.contains(Features::CUBE_TEXTURES_ARRAY) { + if version < Version::Embedded(310) || version < Version::Desktop(130) { + return Err(Error::Custom(format!( + "Version {} doesn't support cube map array textures", + version + ))); + } + + if version.is_es() { + // https://www.khronos.org/registry/OpenGL/extensions/EXT/EXT_texture_cube_map_array.txt + writeln!(out, "#extension GL_EXT_texture_cube_map_array : require")?; + } else if version < Version::Desktop(400) { + // https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_texture_cube_map_array.txt + writeln!(out, "#extension GL_ARB_texture_cube_map_array : require")?; + } + } + + if self.0.contains(Features::MULTISAMPLED_TEXTURES) { + if version < Version::Embedded(300) { + return Err(Error::Custom(format!( + "Version {} doesn't support multi sampled textures", + version + ))); + } + } + + if self.0.contains(Features::MULTISAMPLED_TEXTURE_ARRAYS) { + if version < Version::Embedded(310) { + return Err(Error::Custom(format!( + "Version {} doesn't support multi sampled texture arrays", + version + ))); + } + + if version.is_es() { + // https://www.khronos.org/registry/OpenGL/extensions/OES/OES_texture_storage_multisample_2d_array.txt + writeln!( + out, + "#extension GL_OES_texture_storage_multisample_2d_array : require" + )?; + } + } + + if self.0.contains(Features::ARRAY_OF_ARRAYS) { + if version < Version::Embedded(310) || version < Version::Desktop(120) { + return Err(Error::Custom(format!( + "Version {} doesn't arrays of arrays", + version + ))); + } + + if version < Version::Desktop(430) { + // https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_arrays_of_arrays.txt + writeln!(out, "#extension ARB_arrays_of_arrays : require")?; + } + } + + if self.0.contains(Features::IMAGE_LOAD_STORE) { + if version < Version::Embedded(310) || version < Version::Desktop(130) { + return Err(Error::Custom(format!( + "Version {} doesn't support images load/stores", + version + ))); + } + + if self.0.contains(Features::FULL_IMAGE_FORMATS) && version.is_es() { + // https://www.khronos.org/registry/OpenGL/extensions/NV/NV_image_formats.txt + writeln!(out, "#extension GL_NV_image_formats : require")?; + } + + if version < Version::Desktop(420) { + // https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_shader_image_load_store.txt + writeln!(out, "#extension GL_ARB_shader_image_load_store : require")?; + } + } + + if self.0.contains(Features::CONSERVATIVE_DEPTH) { + if version < Version::Embedded(300) || version < Version::Desktop(130) { + return Err(Error::Custom(format!( + "Version {} doesn't support conservative depth", + version + ))); + } + + if version.is_es() { + // https://www.khronos.org/registry/OpenGL/extensions/EXT/EXT_conservative_depth.txt + writeln!(out, "#extension GL_EXT_conservative_depth : require")?; + } + + if version < Version::Desktop(420) { + // https://www.khronos.org/registry/OpenGL/extensions/ARB/ARB_conservative_depth.txt + writeln!(out, "#extension GL_ARB_conservative_depth : require")?; + } + } + + if self.0.contains(Features::TEXTURE_1D) { + if version.is_es() { + return Err(Error::Custom(format!( + "Version {} doesn't support 1d textures", + version + ))); + } + } + + Ok(()) + } +} + +enum FunctionType { + Function(Handle<Function>), + EntryPoint(crate::proc::EntryPointIndex), +} + +struct FunctionCtx<'a, 'b> { + func: FunctionType, + expressions: &'a Arena<Expression>, + typifier: &'b Typifier, +} + +impl<'a, 'b> FunctionCtx<'a, 'b> { + fn name_key(&self, local: Handle<LocalVariable>) -> NameKey { + match self.func { + FunctionType::Function(handle) => NameKey::FunctionLocal(handle, local), + FunctionType::EntryPoint(idx) => NameKey::EntryPointLocal(idx, local), + } + } + + fn get_arg<'c>(&self, arg: u32, names: &'c FastHashMap<NameKey, String>) -> &'c str { + match self.func { + FunctionType::Function(handle) => &names[&NameKey::FunctionArgument(handle, arg)], + FunctionType::EntryPoint(_) => unreachable!(), + } + } +} + +/// Helper structure that generates a number +#[derive(Default)] +struct IdGenerator(u32); + +impl IdGenerator { + fn generate(&mut self) -> u32 { + let ret = self.0; + self.0 += 1; + ret + } +} + +/// Main structure of the glsl backend responsible for all code generation +pub struct Writer<'a, W> { + // Inputs + module: &'a Module, + out: W, + options: &'a Options, + + // Internal State + features: FeaturesManager, + names: FastHashMap<NameKey, String>, + entry_point: &'a crate::EntryPoint, + entry_point_idx: crate::proc::EntryPointIndex, + call_graph: CallGraph, + + /// Used to generate a unique number for blocks + block_id: IdGenerator, +} + +impl<'a, W: Write> Writer<'a, W> { + pub fn new(out: W, module: &'a Module, options: &'a Options) -> Result<Self, Error> { + if !options.version.is_supported() { + return Err(Error::Custom(format!( + "Version not supported {}", + options.version + ))); + } + + let (ep_idx, ep) = module + .entry_points + .iter() + .enumerate() + .find_map(|(i, (key, entry_point))| { + Some((i as u16, entry_point)).filter(|_| &options.entry_point == key) + }) + .ok_or_else(|| Error::Custom(String::from("Entry point not found")))?; + + let mut names = FastHashMap::default(); + Namer::process(module, keywords::RESERVED_KEYWORDS, &mut names); + + let call_graph = CallGraphBuilder { + functions: &module.functions, + } + .process(&ep.function); + + let mut this = Self { + module, + out, + options, + + features: FeaturesManager::new(), + names, + entry_point: ep, + entry_point_idx: ep_idx, + call_graph, + + block_id: IdGenerator::default(), + }; + + this.collect_required_features()?; + + Ok(this) + } + + fn collect_required_features(&mut self) -> Result<(), Error> { + let stage = self.options.entry_point.0; + + if let Some(depth_test) = self.entry_point.early_depth_test { + self.features.request(Features::IMAGE_LOAD_STORE); + + if depth_test.conservative.is_some() { + self.features.request(Features::CONSERVATIVE_DEPTH); + } + } + + if let ShaderStage::Compute = stage { + self.features.request(Features::COMPUTE_SHADER) + } + + for (_, ty) in self.module.types.iter() { + match ty.inner { + TypeInner::Scalar { kind, width } => self.scalar_required_features(kind, width), + TypeInner::Vector { kind, width, .. } => self.scalar_required_features(kind, width), + TypeInner::Matrix { .. } => self.scalar_required_features(ScalarKind::Float, 8), + TypeInner::Array { base, .. } => { + if let TypeInner::Array { .. } = self.module.types[base].inner { + self.features.request(Features::ARRAY_OF_ARRAYS) + } + } + TypeInner::Image { + dim, + arrayed, + class, + } => { + if arrayed && dim == crate::ImageDimension::Cube { + self.features.request(Features::CUBE_TEXTURES_ARRAY) + } else if dim == crate::ImageDimension::D1 { + self.features.request(Features::TEXTURE_1D) + } + + match class { + ImageClass::Sampled { multi: true, .. } => { + self.features.request(Features::MULTISAMPLED_TEXTURES); + if arrayed { + self.features.request(Features::MULTISAMPLED_TEXTURE_ARRAYS); + } + } + ImageClass::Storage(format) => match format { + StorageFormat::R8Unorm + | StorageFormat::R8Snorm + | StorageFormat::R8Uint + | StorageFormat::R8Sint + | StorageFormat::R16Uint + | StorageFormat::R16Sint + | StorageFormat::R16Float + | StorageFormat::Rg8Unorm + | StorageFormat::Rg8Snorm + | StorageFormat::Rg8Uint + | StorageFormat::Rg8Sint + | StorageFormat::Rg16Uint + | StorageFormat::Rg16Sint + | StorageFormat::Rg16Float + | StorageFormat::Rgb10a2Unorm + | StorageFormat::Rg11b10Float + | StorageFormat::Rg32Uint + | StorageFormat::Rg32Sint + | StorageFormat::Rg32Float => { + self.features.request(Features::FULL_IMAGE_FORMATS) + } + _ => {} + }, + _ => {} + } + } + _ => {} + } + } + + for (_, global) in self.module.global_variables.iter() { + match global.class { + StorageClass::WorkGroup => self.features.request(Features::COMPUTE_SHADER), + StorageClass::Storage => self.features.request(Features::BUFFER_STORAGE), + StorageClass::PushConstant => self.features.request(Features::PUSH_CONSTANT), + _ => {} + } + } + + Ok(()) + } + + fn scalar_required_features(&mut self, kind: ScalarKind, width: crate::Bytes) { + if kind == ScalarKind::Float && width == 8 { + self.features.request(Features::DOUBLE_TYPE); + } + } + + pub fn write(&mut self) -> Result<FastHashMap<String, TextureMapping>, Error> { + let es = self.options.version.is_es(); + + writeln!(self.out, "#version {}", self.options.version)?; + self.features.write(self.options.version, &mut self.out)?; + writeln!(self.out)?; + + if es { + writeln!(self.out, "precision highp float;\n")?; + } + + if let Some(depth_test) = self.entry_point.early_depth_test { + writeln!(self.out, "layout(early_fragment_tests) in;\n")?; + + if let Some(conservative) = depth_test.conservative { + writeln!( + self.out, + "layout (depth_{}) out float gl_FragDepth;\n", + match conservative { + ConservativeDepth::GreaterEqual => "greater", + ConservativeDepth::LessEqual => "less", + ConservativeDepth::Unchanged => "unchanged", + } + )?; + } + } + + for (handle, ty) in self.module.types.iter() { + if let TypeInner::Struct { ref members } = ty.inner { + self.write_struct(handle, members)? + } + } + + writeln!(self.out)?; + + let texture_mappings = self.collect_texture_mapping( + self.call_graph + .raw_nodes() + .iter() + .map(|node| &self.module.functions[node.weight]) + .chain(std::iter::once(&self.entry_point.function)), + )?; + + for (handle, global) in self + .module + .global_variables + .iter() + .zip(&self.entry_point.function.global_usage) + .filter_map(|(global, usage)| Some(global).filter(|_| !usage.is_empty())) + { + if let Some(crate::Binding::BuiltIn(_)) = global.binding { + continue; + } + + match self.module.types[global.ty].inner { + TypeInner::Image { + dim, + arrayed, + class, + } => { + if let TypeInner::Image { + class: ImageClass::Storage(format), + .. + } = self.module.types[global.ty].inner + { + write!(self.out, "layout({}) ", glsl_storage_format(format))?; + } + + if global.storage_access == StorageAccess::LOAD { + write!(self.out, "readonly ")?; + } else if global.storage_access == StorageAccess::STORE { + write!(self.out, "writeonly ")?; + } + + write!(self.out, "uniform ")?; + + self.write_image_type(dim, arrayed, class)?; + + writeln!( + self.out, + " {};", + self.names[&NameKey::GlobalVariable(handle)] + )? + } + TypeInner::Sampler { .. } => continue, + _ => self.write_global(handle, global)?, + } + } + + writeln!(self.out)?; + + // Sort the graph topologically so that functions calls are valid + // It's impossible for this to panic because the IR forbids cycles + let functions = petgraph::algo::toposort(&self.call_graph, None).unwrap(); + + for node in functions { + let handle = self.call_graph[node]; + let name = self.names[&NameKey::Function(handle)].clone(); + self.write_function( + FunctionType::Function(handle), + &self.module.functions[handle], + name, + )?; + } + + self.write_function( + FunctionType::EntryPoint(self.entry_point_idx), + &self.entry_point.function, + "main", + )?; + + Ok(texture_mappings) + } + + fn write_global( + &mut self, + handle: Handle<GlobalVariable>, + global: &GlobalVariable, + ) -> Result<(), Error> { + if global.storage_access == StorageAccess::LOAD { + write!(self.out, "readonly ")?; + } else if global.storage_access == StorageAccess::STORE { + write!(self.out, "writeonly ")?; + } + + if let Some(interpolation) = global.interpolation { + match (self.options.entry_point.0, global.class) { + (ShaderStage::Fragment, StorageClass::Input) + | (ShaderStage::Vertex, StorageClass::Output) => { + write!(self.out, "{} ", glsl_interpolation(interpolation)?)?; + } + _ => (), + }; + } + + let block = match global.class { + StorageClass::Storage | StorageClass::Uniform => { + let block_name = self.names[&NameKey::Type(global.ty)].clone(); + + Some(block_name) + } + _ => None, + }; + + write!(self.out, "{} ", glsl_storage_class(global.class))?; + + self.write_type(global.ty, block)?; + + let name = &self.names[&NameKey::GlobalVariable(handle)]; + writeln!(self.out, " {};", name)?; + + Ok(()) + } + + fn write_function<N: AsRef<str>>( + &mut self, + ty: FunctionType, + func: &Function, + name: N, + ) -> Result<(), Error> { + let mut typifier = Typifier::new(); + + typifier.resolve_all( + &func.expressions, + &self.module.types, + &ResolveContext { + constants: &self.module.constants, + global_vars: &self.module.global_variables, + local_vars: &func.local_variables, + functions: &self.module.functions, + arguments: &func.arguments, + }, + )?; + + let ctx = FunctionCtx { + func: ty, + expressions: &func.expressions, + typifier: &typifier, + }; + + self.write_fn_header(name.as_ref(), func, &ctx)?; + writeln!(self.out, " {{",)?; + + for (handle, local) in func.local_variables.iter() { + write!(self.out, "\t")?; + self.write_type(local.ty, None)?; + + write!(self.out, " {}", self.names[&ctx.name_key(handle)])?; + + if let Some(init) = local.init { + write!(self.out, " = ",)?; + + self.write_constant(&self.module.constants[init])?; + } + + writeln!(self.out, ";")? + } + + writeln!(self.out)?; + + for sta in func.body.iter() { + self.write_stmt(sta, &ctx, 1)?; + } + + Ok(writeln!(self.out, "}}")?) + } + + fn write_slice<T, F: FnMut(&mut Self, u32, &T) -> Result<(), Error>>( + &mut self, + data: &[T], + mut f: F, + ) -> Result<(), Error> { + for (i, item) in data.iter().enumerate() { + f(self, i as u32, item)?; + + if i != data.len().saturating_sub(1) { + write!(self.out, ",")?; + } + } + + Ok(()) + } + + fn write_fn_header( + &mut self, + name: &str, + func: &Function, + ctx: &FunctionCtx<'_, '_>, + ) -> Result<(), Error> { + if let Some(ty) = func.return_type { + self.write_type(ty, None)?; + } else { + write!(self.out, "void")?; + } + + write!(self.out, " {}(", name)?; + + self.write_slice(&func.arguments, |this, i, arg| { + this.write_type(arg.ty, None)?; + + let name = ctx.get_arg(i, &this.names); + + Ok(write!(this.out, " {}", name)?) + })?; + + write!(self.out, ")")?; + + Ok(()) + } + + fn write_type(&mut self, ty: Handle<Type>, block: Option<String>) -> Result<(), Error> { + match self.module.types[ty].inner { + TypeInner::Scalar { kind, width } => { + write!(self.out, "{}", glsl_scalar(kind, width)?.full)? + } + TypeInner::Vector { size, kind, width } => write!( + self.out, + "{}vec{}", + glsl_scalar(kind, width)?.prefix, + size as u8 + )?, + TypeInner::Matrix { + columns, + rows, + width, + } => write!( + self.out, + "{}mat{}x{}", + glsl_scalar(ScalarKind::Float, width)?.prefix, + columns as u8, + rows as u8 + )?, + TypeInner::Pointer { base, .. } => self.write_type(base, None)?, + TypeInner::Array { base, size, .. } => { + self.write_type(base, None)?; + + write!(self.out, "[")?; + self.write_array_size(size)?; + write!(self.out, "]")? + } + TypeInner::Struct { ref members } => { + if let Some(name) = block { + writeln!(self.out, "{}_block_{} {{", name, self.block_id.generate())?; + + for (idx, member) in members.iter().enumerate() { + self.write_type(member.ty, None)?; + + writeln!( + self.out, + " {};", + &self.names[&NameKey::StructMember(ty, idx as u32)] + )?; + } + + write!(self.out, "}}")? + } else { + write!(self.out, "{}", &self.names[&NameKey::Type(ty)])? + } + } + _ => unreachable!(), + } + + Ok(()) + } + + fn write_image_type( + &mut self, + dim: crate::ImageDimension, + arrayed: bool, + class: ImageClass, + ) -> Result<(), Error> { + let (base, kind, ms, comparison) = match class { + ImageClass::Sampled { kind, multi: true } => ("sampler", kind, "MS", ""), + ImageClass::Sampled { kind, multi: false } => ("sampler", kind, "", ""), + ImageClass::Depth => ("sampler", crate::ScalarKind::Float, "", "Shadow"), + ImageClass::Storage(format) => ("image", format.into(), "", ""), + }; + + Ok(write!( + self.out, + "{}{}{}{}{}{}", + glsl_scalar(kind, 4)?.prefix, + base, + ImageDimension(dim), + ms, + if arrayed { "Array" } else { "" }, + comparison + )?) + } + + fn write_array_size(&mut self, size: ArraySize) -> Result<(), Error> { + match size { + ArraySize::Constant(const_handle) => match self.module.constants[const_handle].inner { + ConstantInner::Uint(size) => write!(self.out, "{}", size)?, + _ => unreachable!(), + }, + ArraySize::Dynamic => (), + } + + Ok(()) + } + + fn collect_texture_mapping( + &self, + functions: impl Iterator<Item = &'a Function>, + ) -> Result<FastHashMap<String, TextureMapping>, Error> { + let mut mappings = FastHashMap::default(); + + for func in functions { + let mut interface = Interface { + expressions: &func.expressions, + local_variables: &func.local_variables, + visitor: TextureMappingVisitor { + names: &self.names, + expressions: &func.expressions, + map: &mut mappings, + error: None, + }, + }; + interface.traverse(&func.body); + + if let Some(error) = interface.visitor.error { + return Err(error); + } + } + + Ok(mappings) + } + + fn write_struct( + &mut self, + handle: Handle<Type>, + members: &[StructMember], + ) -> Result<(), Error> { + writeln!(self.out, "struct {} {{", self.names[&NameKey::Type(handle)])?; + + for (idx, member) in members.iter().enumerate() { + write!(self.out, "\t")?; + self.write_type(member.ty, None)?; + writeln!( + self.out, + " {};", + self.names[&NameKey::StructMember(handle, idx as u32)] + )?; + } + + writeln!(self.out, "}};")?; + Ok(()) + } + + fn write_stmt( + &mut self, + sta: &Statement, + ctx: &FunctionCtx<'_, '_>, + indent: usize, + ) -> Result<(), Error> { + write!(self.out, "{}", "\t".repeat(indent))?; + + match sta { + Statement::Block(block) => { + writeln!(self.out, "{{")?; + for sta in block.iter() { + self.write_stmt(sta, ctx, indent + 1)? + } + writeln!(self.out, "{}}}", "\t".repeat(indent))? + } + Statement::If { + condition, + accept, + reject, + } => { + write!(self.out, "if(")?; + self.write_expr(*condition, ctx)?; + writeln!(self.out, ") {{")?; + + for sta in accept { + self.write_stmt(sta, ctx, indent + 1)?; + } + + if !reject.is_empty() { + writeln!(self.out, "{}}} else {{", "\t".repeat(indent))?; + + for sta in reject { + self.write_stmt(sta, ctx, indent + 1)?; + } + } + + writeln!(self.out, "{}}}", "\t".repeat(indent))? + } + Statement::Switch { + selector, + cases, + default, + } => { + write!(self.out, "switch(")?; + self.write_expr(*selector, ctx)?; + writeln!(self.out, ") {{")?; + + for (label, (block, fallthrough)) in cases { + writeln!(self.out, "{}case {}:", "\t".repeat(indent + 1), label)?; + + for sta in block { + self.write_stmt(sta, ctx, indent + 2)?; + } + + if fallthrough.is_none() { + writeln!(self.out, "{}break;", "\t".repeat(indent + 2))?; + } + } + + if !default.is_empty() { + writeln!(self.out, "{}default:", "\t".repeat(indent + 1))?; + + for sta in default { + self.write_stmt(sta, ctx, indent + 2)?; + } + } + + writeln!(self.out, "{}}}", "\t".repeat(indent))? + } + Statement::Loop { body, continuing } => { + writeln!(self.out, "while(true) {{")?; + + for sta in body.iter().chain(continuing.iter()) { + self.write_stmt(sta, ctx, indent + 1)?; + } + + writeln!(self.out, "{}}}", "\t".repeat(indent))? + } + Statement::Break => writeln!(self.out, "break;")?, + Statement::Continue => writeln!(self.out, "continue;")?, + Statement::Return { value } => { + write!(self.out, "return")?; + if let Some(expr) = value { + write!(self.out, " ")?; + self.write_expr(*expr, ctx)?; + } + writeln!(self.out, ";")?; + } + Statement::Kill => writeln!(self.out, "discard;")?, + Statement::Store { pointer, value } => { + self.write_expr(*pointer, ctx)?; + write!(self.out, " = ")?; + self.write_expr(*value, ctx)?; + writeln!(self.out, ";")? + } + } + + Ok(()) + } + + fn write_expr( + &mut self, + expr: Handle<Expression>, + ctx: &FunctionCtx<'_, '_>, + ) -> Result<(), Error> { + match ctx.expressions[expr] { + Expression::Access { base, index } => { + self.write_expr(base, ctx)?; + write!(self.out, "[")?; + self.write_expr(index, ctx)?; + write!(self.out, "]")? + } + Expression::AccessIndex { base, index } => { + self.write_expr(base, ctx)?; + + match ctx.typifier.get(base, &self.module.types) { + TypeInner::Vector { .. } + | TypeInner::Matrix { .. } + | TypeInner::Array { .. } => write!(self.out, "[{}]", index)?, + TypeInner::Struct { .. } => { + let ty = ctx.typifier.get_handle(base).unwrap(); + + write!( + self.out, + ".{}", + &self.names[&NameKey::StructMember(ty, index)] + )? + } + ref other => return Err(Error::Custom(format!("Cannot index {:?}", other))), + } + } + Expression::Constant(constant) => { + self.write_constant(&self.module.constants[constant])? + } + Expression::Compose { ty, ref components } => { + match self.module.types[ty].inner { + TypeInner::Vector { .. } + | TypeInner::Matrix { .. } + | TypeInner::Array { .. } + | TypeInner::Struct { .. } => self.write_type(ty, None)?, + _ => unreachable!(), + } + + write!(self.out, "(")?; + self.write_slice(components, |this, _, arg| this.write_expr(*arg, ctx))?; + write!(self.out, ")")? + } + Expression::FunctionArgument(pos) => { + write!(self.out, "{}", ctx.get_arg(pos, &self.names))? + } + Expression::GlobalVariable(handle) => { + if let Some(crate::Binding::BuiltIn(built_in)) = + self.module.global_variables[handle].binding + { + write!(self.out, "{}", glsl_built_in(built_in))? + } else { + write!( + self.out, + "{}", + &self.names[&NameKey::GlobalVariable(handle)] + )? + } + } + Expression::LocalVariable(handle) => { + write!(self.out, "{}", self.names[&ctx.name_key(handle)])? + } + Expression::Load { pointer } => self.write_expr(pointer, ctx)?, + Expression::ImageSample { + image, + coordinate, + level, + depth_ref, + .. + } => { + //TODO: handle MS + write!( + self.out, + "{}(", + match level { + crate::SampleLevel::Auto | crate::SampleLevel::Bias(_) => "texture", + crate::SampleLevel::Zero | crate::SampleLevel::Exact(_) => "textureLod", + } + )?; + self.write_expr(image, ctx)?; + write!(self.out, ", ")?; + + let size = match *ctx.typifier.get(coordinate, &self.module.types) { + TypeInner::Vector { size, .. } => size, + ref other => { + return Err(Error::Custom(format!( + "Cannot sample with coordinates of type {:?}", + other + ))) + } + }; + + if let Some(depth_ref) = depth_ref { + write!(self.out, "vec{}(", size as u8 + 1)?; + self.write_expr(coordinate, ctx)?; + write!(self.out, ", ")?; + self.write_expr(depth_ref, ctx)?; + write!(self.out, ")")? + } else { + self.write_expr(coordinate, ctx)? + } + + match level { + crate::SampleLevel::Auto => (), + crate::SampleLevel::Zero => write!(self.out, ", 0")?, + crate::SampleLevel::Exact(expr) | crate::SampleLevel::Bias(expr) => { + write!(self.out, ", ")?; + self.write_expr(expr, ctx)?; + } + } + + write!(self.out, ")")? + } + Expression::ImageLoad { + image, + coordinate, + index, + } => { + let class = match ctx.typifier.get(image, &self.module.types) { + TypeInner::Image { class, .. } => class, + _ => unreachable!(), + }; + + match class { + ImageClass::Sampled { .. } => write!(self.out, "texelFetch(")?, + ImageClass::Storage(_) => write!(self.out, "imageLoad(")?, + ImageClass::Depth => todo!(), + } + + self.write_expr(image, ctx)?; + write!(self.out, ", ")?; + self.write_expr(coordinate, ctx)?; + + match class { + ImageClass::Sampled { .. } => { + write!(self.out, ", ")?; + self.write_expr(index.unwrap(), ctx)?; + write!(self.out, ")")? + } + ImageClass::Storage(_) => write!(self.out, ")")?, + ImageClass::Depth => todo!(), + } + } + Expression::Unary { op, expr } => { + write!( + self.out, + "({} ", + match op { + UnaryOperator::Negate => "-", + UnaryOperator::Not => match *ctx.typifier.get(expr, &self.module.types) { + TypeInner::Scalar { + kind: ScalarKind::Sint, + .. + } => "~", + TypeInner::Scalar { + kind: ScalarKind::Uint, + .. + } => "~", + TypeInner::Scalar { + kind: ScalarKind::Bool, + .. + } => "!", + ref other => + return Err(Error::Custom(format!( + "Cannot apply not to type {:?}", + other + ))), + }, + } + )?; + + self.write_expr(expr, ctx)?; + + write!(self.out, ")")? + } + Expression::Binary { op, left, right } => { + write!(self.out, "(")?; + self.write_expr(left, ctx)?; + + write!( + self.out, + " {} ", + match op { + BinaryOperator::Add => "+", + BinaryOperator::Subtract => "-", + BinaryOperator::Multiply => "*", + BinaryOperator::Divide => "/", + BinaryOperator::Modulo => "%", + BinaryOperator::Equal => "==", + BinaryOperator::NotEqual => "!=", + BinaryOperator::Less => "<", + BinaryOperator::LessEqual => "<=", + BinaryOperator::Greater => ">", + BinaryOperator::GreaterEqual => ">=", + BinaryOperator::And => "&", + BinaryOperator::ExclusiveOr => "^", + BinaryOperator::InclusiveOr => "|", + BinaryOperator::LogicalAnd => "&&", + BinaryOperator::LogicalOr => "||", + BinaryOperator::ShiftLeft => "<<", + BinaryOperator::ShiftRight => ">>", + } + )?; + + self.write_expr(right, ctx)?; + + write!(self.out, ")")? + } + Expression::Select { + condition, + accept, + reject, + } => { + write!(self.out, "(")?; + self.write_expr(condition, ctx)?; + write!(self.out, " ? ")?; + self.write_expr(accept, ctx)?; + write!(self.out, " : ")?; + self.write_expr(reject, ctx)?; + write!(self.out, ")")? + } + Expression::Intrinsic { fun, argument } => { + write!( + self.out, + "{}(", + match fun { + IntrinsicFunction::IsFinite => "!isinf", + IntrinsicFunction::IsInf => "isinf", + IntrinsicFunction::IsNan => "isnan", + IntrinsicFunction::IsNormal => "!isnan", + IntrinsicFunction::All => "all", + IntrinsicFunction::Any => "any", + } + )?; + + self.write_expr(argument, ctx)?; + + write!(self.out, ")")? + } + Expression::Transpose(matrix) => { + write!(self.out, "transpose(")?; + self.write_expr(matrix, ctx)?; + write!(self.out, ")")? + } + Expression::DotProduct(left, right) => { + write!(self.out, "dot(")?; + self.write_expr(left, ctx)?; + write!(self.out, ", ")?; + self.write_expr(right, ctx)?; + write!(self.out, ")")? + } + Expression::CrossProduct(left, right) => { + write!(self.out, "cross(")?; + self.write_expr(left, ctx)?; + write!(self.out, ", ")?; + self.write_expr(right, ctx)?; + write!(self.out, ")")? + } + Expression::As { + expr, + kind, + convert, + } => { + if convert { + self.write_type(ctx.typifier.get_handle(expr).unwrap(), None)?; + } else { + let source_kind = match *ctx.typifier.get(expr, &self.module.types) { + TypeInner::Scalar { + kind: source_kind, .. + } => source_kind, + TypeInner::Vector { + kind: source_kind, .. + } => source_kind, + _ => unreachable!(), + }; + + write!( + self.out, + "{}", + match (source_kind, kind) { + (ScalarKind::Float, ScalarKind::Sint) => "floatBitsToInt", + (ScalarKind::Float, ScalarKind::Uint) => "floatBitsToUInt", + (ScalarKind::Sint, ScalarKind::Float) => "intBitsToFloat", + (ScalarKind::Uint, ScalarKind::Float) => "uintBitsToFloat", + _ => { + return Err(Error::Custom(format!( + "Cannot bitcast {:?} to {:?}", + source_kind, kind + ))); + } + } + )?; + } + + write!(self.out, "(")?; + self.write_expr(expr, ctx)?; + write!(self.out, ")")? + } + Expression::Derivative { axis, expr } => { + write!( + self.out, + "{}(", + match axis { + DerivativeAxis::X => "dFdx", + DerivativeAxis::Y => "dFdy", + DerivativeAxis::Width => "fwidth", + } + )?; + self.write_expr(expr, ctx)?; + write!(self.out, ")")? + } + Expression::Call { + origin: FunctionOrigin::Local(ref function), + ref arguments, + } => { + write!(self.out, "{}(", &self.names[&NameKey::Function(*function)])?; + self.write_slice(arguments, |this, _, arg| this.write_expr(*arg, ctx))?; + write!(self.out, ")")? + } + Expression::Call { + origin: crate::FunctionOrigin::External(ref name), + ref arguments, + } => match name.as_str() { + "cos" | "normalize" | "sin" | "length" | "abs" | "floor" | "inverse" + | "distance" | "dot" | "min" | "max" | "reflect" | "pow" | "step" | "cross" + | "fclamp" | "clamp" | "mix" | "smoothstep" => { + let name = match name.as_str() { + "fclamp" => "clamp", + name => name, + }; + + write!(self.out, "{}(", name)?; + self.write_slice(arguments, |this, _, arg| this.write_expr(*arg, ctx))?; + write!(self.out, ")")? + } + "atan2" => { + write!(self.out, "atan(")?; + self.write_expr(arguments[1], ctx)?; + write!(self.out, ", ")?; + self.write_expr(arguments[0], ctx)?; + write!(self.out, ")")? + } + other => { + return Err(Error::Custom(format!( + "Unsupported function call {}", + other + ))) + } + }, + Expression::ArrayLength(expr) => { + write!(self.out, "uint(")?; + self.write_expr(expr, ctx)?; + write!(self.out, ".length())")? + } + } + + Ok(()) + } + + fn write_constant(&mut self, constant: &Constant) -> Result<(), Error> { + match constant.inner { + ConstantInner::Sint(int) => write!(self.out, "{}", int)?, + ConstantInner::Uint(int) => write!(self.out, "{}u", int)?, + ConstantInner::Float(float) => write!(self.out, "{:?}", float)?, + ConstantInner::Bool(boolean) => write!(self.out, "{}", boolean)?, + ConstantInner::Composite(ref components) => { + self.write_type(constant.ty, None)?; + write!(self.out, "(")?; + self.write_slice(components, |this, _, arg| { + this.write_constant(&this.module.constants[*arg]) + })?; + write!(self.out, ")")? + } + } + + Ok(()) + } +} + +struct ScalarString<'a> { + prefix: &'a str, + full: &'a str, +} + +fn glsl_scalar(kind: ScalarKind, width: crate::Bytes) -> Result<ScalarString<'static>, Error> { + Ok(match kind { + ScalarKind::Sint => ScalarString { + prefix: "i", + full: "int", + }, + ScalarKind::Uint => ScalarString { + prefix: "u", + full: "uint", + }, + ScalarKind::Float => match width { + 4 => ScalarString { + prefix: "", + full: "float", + }, + 8 => ScalarString { + prefix: "d", + full: "double", + }, + _ => { + return Err(Error::Custom(format!( + "Cannot build float of width {}", + width + ))) + } + }, + ScalarKind::Bool => ScalarString { + prefix: "b", + full: "bool", + }, + }) +} + +fn glsl_built_in(built_in: BuiltIn) -> &'static str { + match built_in { + BuiltIn::Position => "gl_Position", + BuiltIn::GlobalInvocationId => "gl_GlobalInvocationID", + BuiltIn::BaseInstance => "gl_BaseInstance", + BuiltIn::BaseVertex => "gl_BaseVertex", + BuiltIn::ClipDistance => "gl_ClipDistance", + BuiltIn::InstanceIndex => "gl_InstanceIndex", + BuiltIn::VertexIndex => "gl_VertexIndex", + BuiltIn::PointSize => "gl_PointSize", + BuiltIn::FragCoord => "gl_FragCoord", + BuiltIn::FrontFacing => "gl_FrontFacing", + BuiltIn::SampleIndex => "gl_SampleID", + BuiltIn::FragDepth => "gl_FragDepth", + BuiltIn::LocalInvocationId => "gl_LocalInvocationID", + BuiltIn::LocalInvocationIndex => "gl_LocalInvocationIndex", + BuiltIn::WorkGroupId => "gl_WorkGroupID", + } +} + +fn glsl_storage_class(class: StorageClass) -> &'static str { + match class { + StorageClass::Function => "", + StorageClass::Input => "in", + StorageClass::Output => "out", + StorageClass::Private => "", + StorageClass::Storage => "buffer", + StorageClass::Uniform => "uniform", + StorageClass::Handle => "uniform", + StorageClass::WorkGroup => "shared", + StorageClass::PushConstant => "", + } +} + +fn glsl_interpolation(interpolation: Interpolation) -> Result<&'static str, Error> { + Ok(match interpolation { + Interpolation::Perspective => "smooth", + Interpolation::Linear => "noperspective", + Interpolation::Flat => "flat", + Interpolation::Centroid => "centroid", + Interpolation::Sample => "sample", + Interpolation::Patch => { + return Err(Error::Custom( + "patch interpolation qualifier not supported".to_string(), + )) + } + }) +} + +struct ImageDimension(crate::ImageDimension); +impl fmt::Display for ImageDimension { + fn fmt(&self, f: &mut fmt::Formatter<'_>) -> fmt::Result { + write!( + f, + "{}", + match self.0 { + crate::ImageDimension::D1 => "1D", + crate::ImageDimension::D2 => "2D", + crate::ImageDimension::D3 => "3D", + crate::ImageDimension::Cube => "Cube", + } + ) + } +} + +fn glsl_storage_format(format: StorageFormat) -> &'static str { + match format { + StorageFormat::R8Unorm => "r8", + StorageFormat::R8Snorm => "r8_snorm", + StorageFormat::R8Uint => "r8ui", + StorageFormat::R8Sint => "r8i", + StorageFormat::R16Uint => "r16ui", + StorageFormat::R16Sint => "r16i", + StorageFormat::R16Float => "r16f", + StorageFormat::Rg8Unorm => "rg8", + StorageFormat::Rg8Snorm => "rg8_snorm", + StorageFormat::Rg8Uint => "rg8ui", + StorageFormat::Rg8Sint => "rg8i", + StorageFormat::R32Uint => "r32ui", + StorageFormat::R32Sint => "r32i", + StorageFormat::R32Float => "r32f", + StorageFormat::Rg16Uint => "rg16ui", + StorageFormat::Rg16Sint => "rg16i", + StorageFormat::Rg16Float => "rg16f", + StorageFormat::Rgba8Unorm => "rgba8ui", + StorageFormat::Rgba8Snorm => "rgba8_snorm", + StorageFormat::Rgba8Uint => "rgba8ui", + StorageFormat::Rgba8Sint => "rgba8i", + StorageFormat::Rgb10a2Unorm => "rgb10_a2ui", + StorageFormat::Rg11b10Float => "r11f_g11f_b10f", + StorageFormat::Rg32Uint => "rg32ui", + StorageFormat::Rg32Sint => "rg32i", + StorageFormat::Rg32Float => "rg32f", + StorageFormat::Rgba16Uint => "rgba16ui", + StorageFormat::Rgba16Sint => "rgba16i", + StorageFormat::Rgba16Float => "rgba16f", + StorageFormat::Rgba32Uint => "rgba32ui", + StorageFormat::Rgba32Sint => "rgba32i", + StorageFormat::Rgba32Float => "rgba32f", + } +} + +struct TextureMappingVisitor<'a> { + names: &'a FastHashMap<NameKey, String>, + expressions: &'a Arena<Expression>, + map: &'a mut FastHashMap<String, TextureMapping>, + error: Option<Error>, +} + +impl<'a> Visitor for TextureMappingVisitor<'a> { + fn visit_expr(&mut self, expr: &crate::Expression) { + match expr { + Expression::ImageSample { image, sampler, .. } => { + let tex_handle = match self.expressions[*image] { + Expression::GlobalVariable(global) => global, + _ => unreachable!(), + }; + let tex_name = self.names[&NameKey::GlobalVariable(tex_handle)].clone(); + + let sampler_handle = match self.expressions[*sampler] { + Expression::GlobalVariable(global) => global, + _ => unreachable!(), + }; + + let mapping = self.map.entry(tex_name).or_insert(TextureMapping { + texture: tex_handle, + sampler: Some(sampler_handle), + }); + + if mapping.sampler != Some(sampler_handle) { + self.error = Some(Error::Custom(String::from( + "Cannot use texture with two different samplers", + ))); + } + } + Expression::ImageLoad { image, .. } => { + let tex_handle = match self.expressions[*image] { + Expression::GlobalVariable(global) => global, + _ => unreachable!(), + }; + let tex_name = self.names[&NameKey::GlobalVariable(tex_handle)].clone(); + + let mapping = self.map.entry(tex_name).or_insert(TextureMapping { + texture: tex_handle, + sampler: None, + }); + + if mapping.sampler != None { + self.error = Some(Error::Custom(String::from( + "Cannot use texture with two different samplers", + ))); + } + } + _ => {} + } + } +} diff --git a/third_party/rust/naga/src/back/glsl/keywords.rs b/third_party/rust/naga/src/back/glsl/keywords.rs new file mode 100644 index 0000000000..5a2836c189 --- /dev/null +++ b/third_party/rust/naga/src/back/glsl/keywords.rs @@ -0,0 +1,204 @@ +pub const RESERVED_KEYWORDS: &[&str] = &[ + "attribute", + "const", + "uniform", + "varying", + "buffer", + "shared", + "coherent", + "volatile", + "restrict", + "readonly", + "writeonly", + "atomic_uint", + "layout", + "centroid", + "flat", + "smooth", + "noperspective", + "patch", + "sample", + "break", + "continue", + "do", + "for", + "while", + "switch", + "case", + "default", + "if", + "else", + "subroutine", + "in", + "out", + "inout", + "float", + "double", + "int", + "void", + "bool", + "true", + "false", + "invariant", + "precise", + "discard", + "return", + "mat2", + "mat3", + "mat4", + "dmat2", + "dmat3", + "dmat4", + "mat2x2", + "mat2x3", + "mat2x4", + "dmat2x2", + "dmat2x3", + "dmat2x4", + "mat3x2", + "mat3x3", + "mat3x4", + "dmat3x2", + "dmat3x3", + "dmat3x4", + "mat4x2", + "mat4x3", + "mat4x4", + "dmat4x2", + "dmat4x3", + "dmat4x4", + "vec2", + "vec3", + "vec4", + "ivec2", + "ivec3", + "ivec4", + "bvec2", + "bvec3", + "bvec4", + "dvec2", + "dvec3", + "dvec4", + "uint", + "uvec2", + "uvec3", + "uvec4", + "lowp", + "mediump", + "highp", + "precision", + "sampler1D", + "sampler2D", + "sampler3D", + "samplerCube", + "sampler1DShadow", + "sampler2DShadow", + "samplerCubeShadow", + "sampler1DArray", + "sampler2DArray", + "sampler1DArrayShadow", + "sampler2DArrayShadow", + "isampler1D", + "isampler2D", + "isampler3D", + "isamplerCube", + "isampler1DArray", + "isampler2DArray", + "usampler1D", + "usampler2D", + "usampler3D", + "usamplerCube", + "usampler1DArray", + "usampler2DArray", + "sampler2DRect", + "sampler2DRectShadow", + "isampler2D", + "Rect", + "usampler2DRect", + "samplerBuffer", + "isamplerBuffer", + "usamplerBuffer", + "sampler2DMS", + "isampler2DMS", + "usampler2DMS", + "sampler2DMSArray", + "isampler2DMSArray", + "usampler2DMSArray", + "samplerCubeArray", + "samplerCubeArrayShadow", + "isamplerCubeArray", + "usamplerCubeArray", + "image1D", + "iimage1D", + "uimage1D", + "image2D", + "iimage2D", + "uimage2D", + "image3D", + "iimage3D", + "uimage3D", + "image2DRect", + "iimage2DRect", + "uimage2DRect", + "imageCube", + "iimageCube", + "uimageCube", + "imageBuffer", + "iimageBuffer", + "uimageBuffer", + "image1DArray", + "iimage1DArray", + "uimage1DArray", + "image2DArray", + "iimage2DArray", + "uimage2DArray", + "imageCubeArray", + "iimageCubeArray", + "uimageCubeArray", + "image2DMS", + "iimage2DMS", + "uimage2DMS", + "image2DMSArray", + "iimage2DMSArray", + "uimage2DMSArraystruct", + "common", + "partition", + "active", + "asm", + "class", + "union", + "enum", + "typedef", + "template", + "this", + "resource", + "goto", + "inline", + "noinline", + "public", + "static", + "extern", + "external", + "interface", + "long", + "short", + "half", + "fixed", + "unsigned", + "superp", + "input", + "output", + "hvec2", + "hvec3", + "hvec4", + "fvec2", + "fvec3", + "fvec4", + "sampler3DRect", + "filter", + "sizeof", + "cast", + "namespace", + "using", + "main", +]; diff --git a/third_party/rust/naga/src/back/mod.rs b/third_party/rust/naga/src/back/mod.rs new file mode 100644 index 0000000000..bc96dd3496 --- /dev/null +++ b/third_party/rust/naga/src/back/mod.rs @@ -0,0 +1,8 @@ +//! Functions which export shader modules into binary and text formats. + +#[cfg(feature = "glsl-out")] +pub mod glsl; +#[cfg(feature = "msl-out")] +pub mod msl; +#[cfg(feature = "spv-out")] +pub mod spv; diff --git a/third_party/rust/naga/src/back/msl/keywords.rs b/third_party/rust/naga/src/back/msl/keywords.rs new file mode 100644 index 0000000000..cd074ab43f --- /dev/null +++ b/third_party/rust/naga/src/back/msl/keywords.rs @@ -0,0 +1,102 @@ +//TODO: find a complete list +pub const RESERVED: &[&str] = &[ + // control flow + "break", + "if", + "else", + "continue", + "goto", + "do", + "while", + "for", + "switch", + "case", + // types and values + "void", + "unsigned", + "signed", + "bool", + "char", + "int", + "long", + "float", + "double", + "char8_t", + "wchar_t", + "true", + "false", + "nullptr", + "union", + "class", + "struct", + "enum", + // other + "main", + "using", + "decltype", + "sizeof", + "typeof", + "typedef", + "explicit", + "export", + "friend", + "namespace", + "operator", + "public", + "template", + "typename", + "typeid", + "co_await", + "co_return", + "co_yield", + "module", + "import", + "ray_data", + "vec_step", + "visible", + "as_type", + // qualifiers + "mutable", + "static", + "volatile", + "restrict", + "const", + "non-temporal", + "dereferenceable", + "invariant", + // exceptions + "throw", + "try", + "catch", + // operators + "const_cast", + "dynamic_cast", + "reinterpret_cast", + "static_cast", + "new", + "delete", + "and", + "and_eq", + "bitand", + "bitor", + "compl", + "not", + "not_eq", + "or", + "or_eq", + "xor", + "xor_eq", + "compl", + // Metal-specific + "constant", + "device", + "threadgroup", + "threadgroup_imageblock", + "kernel", + "compute", + "vertex", + "fragment", + "read_only", + "write_only", + "read_write", +]; diff --git a/third_party/rust/naga/src/back/msl/mod.rs b/third_party/rust/naga/src/back/msl/mod.rs new file mode 100644 index 0000000000..493e7d0c85 --- /dev/null +++ b/third_party/rust/naga/src/back/msl/mod.rs @@ -0,0 +1,211 @@ +/*! Metal Shading Language (MSL) backend + +## Binding model + +Metal's bindings are flat per resource. Since there isn't an obvious mapping +from SPIR-V's descriptor sets, we require a separate mapping provided in the options. +This mapping may have one or more resource end points for each descriptor set + index +pair. + +## Outputs + +In Metal, built-in shader outputs can not be nested into structures within +the output struct. If there is a structure in the outputs, and it contains any built-ins, +we move them up to the root output structure that we define ourselves. +!*/ + +use crate::{arena::Handle, proc::ResolveError, FastHashMap}; +use std::{ + io::{Error as IoError, Write}, + string::FromUtf8Error, +}; + +mod keywords; +mod writer; + +pub use writer::Writer; + +#[derive(Clone, Debug, Default, PartialEq)] +pub struct BindTarget { + pub buffer: Option<u8>, + pub texture: Option<u8>, + pub sampler: Option<u8>, + pub mutable: bool, +} + +#[derive(Clone, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)] +pub struct BindSource { + pub stage: crate::ShaderStage, + pub group: u32, + pub binding: u32, +} + +pub type BindingMap = FastHashMap<BindSource, BindTarget>; + +enum ResolvedBinding { + BuiltIn(crate::BuiltIn), + Attribute(u32), + Color(u32), + User { prefix: &'static str, index: u32 }, + Resource(BindTarget), +} + +// Note: some of these should be removed in favor of proper IR validation. + +#[derive(Debug)] +pub enum Error { + IO(IoError), + Utf8(FromUtf8Error), + Type(ResolveError), + UnexpectedLocation, + MissingBinding(Handle<crate::GlobalVariable>), + MissingBindTarget(BindSource), + InvalidImageAccess(crate::StorageAccess), + MutabilityViolation(Handle<crate::GlobalVariable>), + BadName(String), + UnexpectedGlobalType(Handle<crate::Type>), + UnimplementedBindTarget(BindTarget), + UnsupportedCompose(Handle<crate::Type>), + UnsupportedBinaryOp(crate::BinaryOperator), + UnexpectedSampleLevel(crate::SampleLevel), + UnsupportedCall(String), + UnsupportedDynamicArrayLength, + UnableToReturnValue(Handle<crate::Expression>), + /// The source IR is not valid. + Validation, +} + +impl From<IoError> for Error { + fn from(e: IoError) -> Self { + Error::IO(e) + } +} + +impl From<FromUtf8Error> for Error { + fn from(e: FromUtf8Error) -> Self { + Error::Utf8(e) + } +} + +impl From<ResolveError> for Error { + fn from(e: ResolveError) -> Self { + Error::Type(e) + } +} + +#[derive(Clone, Copy, Debug)] +enum LocationMode { + VertexInput, + FragmentOutput, + Intermediate, + Uniform, +} + +#[derive(Debug, Default, Clone)] +pub struct Options { + /// (Major, Minor) target version of the Metal Shading Language. + pub lang_version: (u8, u8), + /// Make it possible to link different stages via SPIRV-Cross. + pub spirv_cross_compatibility: bool, + /// Binding model mapping to Metal. + pub binding_map: BindingMap, +} + +impl Options { + fn resolve_binding( + &self, + stage: crate::ShaderStage, + binding: &crate::Binding, + mode: LocationMode, + ) -> Result<ResolvedBinding, Error> { + match *binding { + crate::Binding::BuiltIn(built_in) => Ok(ResolvedBinding::BuiltIn(built_in)), + crate::Binding::Location(index) => match mode { + LocationMode::VertexInput => Ok(ResolvedBinding::Attribute(index)), + LocationMode::FragmentOutput => Ok(ResolvedBinding::Color(index)), + LocationMode::Intermediate => Ok(ResolvedBinding::User { + prefix: if self.spirv_cross_compatibility { + "locn" + } else { + "loc" + }, + index, + }), + LocationMode::Uniform => Err(Error::UnexpectedLocation), + }, + crate::Binding::Resource { group, binding } => { + let source = BindSource { + stage, + group, + binding, + }; + self.binding_map + .get(&source) + .cloned() + .map(ResolvedBinding::Resource) + .ok_or(Error::MissingBindTarget(source)) + } + } + } +} + +impl ResolvedBinding { + fn try_fmt<W: Write>(&self, out: &mut W) -> Result<(), Error> { + match *self { + ResolvedBinding::BuiltIn(built_in) => { + use crate::BuiltIn as Bi; + let name = match built_in { + // vertex + Bi::BaseInstance => "base_instance", + Bi::BaseVertex => "base_vertex", + Bi::ClipDistance => "clip_distance", + Bi::InstanceIndex => "instance_id", + Bi::PointSize => "point_size", + Bi::Position => "position", + Bi::VertexIndex => "vertex_id", + // fragment + Bi::FragCoord => "position", + Bi::FragDepth => "depth(any)", + Bi::FrontFacing => "front_facing", + Bi::SampleIndex => "sample_id", + // compute + Bi::GlobalInvocationId => "thread_position_in_grid", + Bi::LocalInvocationId => "thread_position_in_threadgroup", + Bi::LocalInvocationIndex => "thread_index_in_threadgroup", + Bi::WorkGroupId => "threadgroup_position_in_grid", + }; + Ok(write!(out, "{}", name)?) + } + ResolvedBinding::Attribute(index) => Ok(write!(out, "attribute({})", index)?), + ResolvedBinding::Color(index) => Ok(write!(out, "color({})", index)?), + ResolvedBinding::User { prefix, index } => { + Ok(write!(out, "user({}{})", prefix, index)?) + } + ResolvedBinding::Resource(ref target) => { + if let Some(id) = target.buffer { + Ok(write!(out, "buffer({})", id)?) + } else if let Some(id) = target.texture { + Ok(write!(out, "texture({})", id)?) + } else if let Some(id) = target.sampler { + Ok(write!(out, "sampler({})", id)?) + } else { + Err(Error::UnimplementedBindTarget(target.clone())) + } + } + } + } + + fn try_fmt_decorated<W: Write>(&self, out: &mut W, terminator: &str) -> Result<(), Error> { + write!(out, " [[")?; + self.try_fmt(out)?; + write!(out, "]]")?; + write!(out, "{}", terminator)?; + Ok(()) + } +} + +pub fn write_string(module: &crate::Module, options: &Options) -> Result<String, Error> { + let mut w = writer::Writer::new(Vec::new()); + w.write(module, options)?; + Ok(String::from_utf8(w.finish())?) +} diff --git a/third_party/rust/naga/src/back/msl/writer.rs b/third_party/rust/naga/src/back/msl/writer.rs new file mode 100644 index 0000000000..e2adbea6b7 --- /dev/null +++ b/third_party/rust/naga/src/back/msl/writer.rs @@ -0,0 +1,990 @@ +use super::{keywords::RESERVED, Error, LocationMode, Options, ResolvedBinding}; +use crate::{ + arena::Handle, + proc::{EntryPointIndex, NameKey, Namer, ResolveContext, Typifier}, + FastHashMap, +}; +use std::{ + fmt::{Display, Error as FmtError, Formatter}, + io::Write, +}; + +struct Level(usize); +impl Level { + fn next(&self) -> Self { + Level(self.0 + 1) + } +} +impl Display for Level { + fn fmt(&self, formatter: &mut Formatter<'_>) -> Result<(), FmtError> { + (0..self.0).map(|_| formatter.write_str("\t")).collect() + } +} + +struct TypedGlobalVariable<'a> { + module: &'a crate::Module, + names: &'a FastHashMap<NameKey, String>, + handle: Handle<crate::GlobalVariable>, + usage: crate::GlobalUse, +} + +impl<'a> TypedGlobalVariable<'a> { + fn try_fmt<W: Write>(&self, out: &mut W) -> Result<(), Error> { + let var = &self.module.global_variables[self.handle]; + let name = &self.names[&NameKey::GlobalVariable(self.handle)]; + let ty = &self.module.types[var.ty]; + let ty_name = &self.names[&NameKey::Type(var.ty)]; + + let (space_qualifier, reference) = match ty.inner { + crate::TypeInner::Struct { .. } => match var.class { + crate::StorageClass::Uniform | crate::StorageClass::Storage => { + let space = if self.usage.contains(crate::GlobalUse::STORE) { + "device " + } else { + "constant " + }; + (space, "&") + } + _ => ("", ""), + }, + _ => ("", ""), + }; + Ok(write!( + out, + "{}{}{} {}", + space_qualifier, ty_name, reference, name + )?) + } +} + +pub struct Writer<W> { + out: W, + names: FastHashMap<NameKey, String>, + typifier: Typifier, +} + +fn scalar_kind_string(kind: crate::ScalarKind) -> &'static str { + match kind { + crate::ScalarKind::Float => "float", + crate::ScalarKind::Sint => "int", + crate::ScalarKind::Uint => "uint", + crate::ScalarKind::Bool => "bool", + } +} + +fn vector_size_string(size: crate::VectorSize) -> &'static str { + match size { + crate::VectorSize::Bi => "2", + crate::VectorSize::Tri => "3", + crate::VectorSize::Quad => "4", + } +} + +const OUTPUT_STRUCT_NAME: &str = "output"; +const LOCATION_INPUT_STRUCT_NAME: &str = "input"; +const COMPONENTS: &[char] = &['x', 'y', 'z', 'w']; + +fn separate(is_last: bool) -> &'static str { + if is_last { + "" + } else { + "," + } +} + +enum FunctionOrigin { + Handle(Handle<crate::Function>), + EntryPoint(EntryPointIndex), +} + +struct ExpressionContext<'a> { + function: &'a crate::Function, + origin: FunctionOrigin, + module: &'a crate::Module, +} + +impl<W: Write> Writer<W> { + /// Creates a new `Writer` instance. + pub fn new(out: W) -> Self { + Writer { + out, + names: FastHashMap::default(), + typifier: Typifier::new(), + } + } + + /// Finishes writing and returns the output. + pub fn finish(self) -> W { + self.out + } + + fn put_call( + &mut self, + name: &str, + parameters: &[Handle<crate::Expression>], + context: &ExpressionContext, + ) -> Result<(), Error> { + if !name.is_empty() { + write!(self.out, "metal::{}", name)?; + } + write!(self.out, "(")?; + for (i, &handle) in parameters.iter().enumerate() { + if i != 0 { + write!(self.out, ", ")?; + } + self.put_expression(handle, context)?; + } + write!(self.out, ")")?; + Ok(()) + } + + fn put_expression( + &mut self, + expr_handle: Handle<crate::Expression>, + context: &ExpressionContext, + ) -> Result<(), Error> { + let expression = &context.function.expressions[expr_handle]; + log::trace!("expression {:?} = {:?}", expr_handle, expression); + match *expression { + crate::Expression::Access { base, index } => { + self.put_expression(base, context)?; + write!(self.out, "[")?; + self.put_expression(index, context)?; + write!(self.out, "]")?; + } + crate::Expression::AccessIndex { base, index } => { + self.put_expression(base, context)?; + let resolved = self.typifier.get(base, &context.module.types); + match *resolved { + crate::TypeInner::Struct { .. } => { + let base_ty = self.typifier.get_handle(base).unwrap(); + let name = &self.names[&NameKey::StructMember(base_ty, index)]; + write!(self.out, ".{}", name)?; + } + crate::TypeInner::Matrix { .. } | crate::TypeInner::Vector { .. } => { + write!(self.out, ".{}", COMPONENTS[index as usize])?; + } + crate::TypeInner::Array { .. } => { + write!(self.out, "[{}]", index)?; + } + _ => { + // unexpected indexing, should fail validation + } + } + } + crate::Expression::Constant(handle) => self.put_constant(handle, context.module)?, + crate::Expression::Compose { ty, ref components } => { + let inner = &context.module.types[ty].inner; + match *inner { + crate::TypeInner::Vector { size, kind, .. } => { + write!( + self.out, + "{}{}", + scalar_kind_string(kind), + vector_size_string(size) + )?; + self.put_call("", components, context)?; + } + crate::TypeInner::Scalar { width: 4, kind } if components.len() == 1 => { + write!(self.out, "{}", scalar_kind_string(kind),)?; + self.put_call("", components, context)?; + } + _ => return Err(Error::UnsupportedCompose(ty)), + } + } + crate::Expression::FunctionArgument(index) => { + let fun_handle = match context.origin { + FunctionOrigin::Handle(handle) => handle, + FunctionOrigin::EntryPoint(_) => unreachable!(), + }; + let name = &self.names[&NameKey::FunctionArgument(fun_handle, index)]; + write!(self.out, "{}", name)?; + } + crate::Expression::GlobalVariable(handle) => { + let var = &context.module.global_variables[handle]; + match var.class { + crate::StorageClass::Output => { + if let crate::TypeInner::Struct { .. } = context.module.types[var.ty].inner + { + return Ok(()); + } + write!(self.out, "{}.", OUTPUT_STRUCT_NAME)?; + } + crate::StorageClass::Input => { + if let Some(crate::Binding::Location(_)) = var.binding { + write!(self.out, "{}.", LOCATION_INPUT_STRUCT_NAME)?; + } + } + _ => {} + } + let name = &self.names[&NameKey::GlobalVariable(handle)]; + write!(self.out, "{}", name)?; + } + crate::Expression::LocalVariable(handle) => { + let name_key = match context.origin { + FunctionOrigin::Handle(fun_handle) => { + NameKey::FunctionLocal(fun_handle, handle) + } + FunctionOrigin::EntryPoint(ep_index) => { + NameKey::EntryPointLocal(ep_index, handle) + } + }; + let name = &self.names[&name_key]; + write!(self.out, "{}", name)?; + } + crate::Expression::Load { pointer } => { + //write!(self.out, "*")?; + self.put_expression(pointer, context)?; + } + crate::Expression::ImageSample { + image, + sampler, + coordinate, + level, + depth_ref, + } => { + let op = match depth_ref { + Some(_) => "sample_compare", + None => "sample", + }; + //TODO: handle arrayed images + self.put_expression(image, context)?; + write!(self.out, ".{}(", op)?; + self.put_expression(sampler, context)?; + write!(self.out, ", ")?; + self.put_expression(coordinate, context)?; + if let Some(dref) = depth_ref { + write!(self.out, ", ")?; + self.put_expression(dref, context)?; + } + match level { + crate::SampleLevel::Auto => {} + crate::SampleLevel::Zero => { + write!(self.out, ", level(0)")?; + } + crate::SampleLevel::Exact(h) => { + write!(self.out, ", level(")?; + self.put_expression(h, context)?; + write!(self.out, ")")?; + } + crate::SampleLevel::Bias(h) => { + write!(self.out, ", bias(")?; + self.put_expression(h, context)?; + write!(self.out, ")")?; + } + } + write!(self.out, ")")?; + } + crate::Expression::ImageLoad { + image, + coordinate, + index, + } => { + //TODO: handle arrayed images + self.put_expression(image, context)?; + write!(self.out, ".read(")?; + self.put_expression(coordinate, context)?; + if let Some(index) = index { + write!(self.out, ", ")?; + self.put_expression(index, context)?; + } + write!(self.out, ")")?; + } + crate::Expression::Unary { op, expr } => { + let op_str = match op { + crate::UnaryOperator::Negate => "-", + crate::UnaryOperator::Not => "!", + }; + write!(self.out, "{}", op_str)?; + self.put_expression(expr, context)?; + } + crate::Expression::Binary { op, left, right } => { + let op_str = match op { + crate::BinaryOperator::Add => "+", + crate::BinaryOperator::Subtract => "-", + crate::BinaryOperator::Multiply => "*", + crate::BinaryOperator::Divide => "/", + crate::BinaryOperator::Modulo => "%", + crate::BinaryOperator::Equal => "==", + crate::BinaryOperator::NotEqual => "!=", + crate::BinaryOperator::Less => "<", + crate::BinaryOperator::LessEqual => "<=", + crate::BinaryOperator::Greater => "==", + crate::BinaryOperator::GreaterEqual => ">=", + crate::BinaryOperator::And => "&", + _ => return Err(Error::UnsupportedBinaryOp(op)), + }; + let kind = self + .typifier + .get(left, &context.module.types) + .scalar_kind() + .ok_or(Error::UnsupportedBinaryOp(op))?; + if op == crate::BinaryOperator::Modulo && kind == crate::ScalarKind::Float { + write!(self.out, "fmod(")?; + self.put_expression(left, context)?; + write!(self.out, ", ")?; + self.put_expression(right, context)?; + write!(self.out, ")")?; + } else { + //write!(self.out, "(")?; + self.put_expression(left, context)?; + write!(self.out, " {} ", op_str)?; + self.put_expression(right, context)?; + //write!(self.out, ")")?; + } + } + crate::Expression::Select { + condition, + accept, + reject, + } => { + write!(self.out, "(")?; + self.put_expression(condition, context)?; + write!(self.out, " ? ")?; + self.put_expression(accept, context)?; + write!(self.out, " : ")?; + self.put_expression(reject, context)?; + write!(self.out, ")")?; + } + crate::Expression::Intrinsic { fun, argument } => { + let op = match fun { + crate::IntrinsicFunction::Any => "any", + crate::IntrinsicFunction::All => "all", + crate::IntrinsicFunction::IsNan => "", + crate::IntrinsicFunction::IsInf => "", + crate::IntrinsicFunction::IsFinite => "", + crate::IntrinsicFunction::IsNormal => "", + }; + self.put_call(op, &[argument], context)?; + } + crate::Expression::Transpose(expr) => { + self.put_call("transpose", &[expr], context)?; + } + crate::Expression::DotProduct(a, b) => { + self.put_call("dot", &[a, b], context)?; + } + crate::Expression::CrossProduct(a, b) => { + self.put_call("cross", &[a, b], context)?; + } + crate::Expression::As { + expr, + kind, + convert, + } => { + let scalar = scalar_kind_string(kind); + let size = match *self.typifier.get(expr, &context.module.types) { + crate::TypeInner::Scalar { .. } => "", + crate::TypeInner::Vector { size, .. } => vector_size_string(size), + _ => return Err(Error::Validation), + }; + let op = if convert { "static_cast" } else { "as_type" }; + write!(self.out, "{}<{}{}>(", op, scalar, size)?; + self.put_expression(expr, context)?; + write!(self.out, ")")?; + } + crate::Expression::Derivative { axis, expr } => { + let op = match axis { + crate::DerivativeAxis::X => "dfdx", + crate::DerivativeAxis::Y => "dfdy", + crate::DerivativeAxis::Width => "fwidth", + }; + self.put_call(op, &[expr], context)?; + } + crate::Expression::Call { + origin: crate::FunctionOrigin::Local(handle), + ref arguments, + } => { + let name = &self.names[&NameKey::Function(handle)]; + write!(self.out, "{}", name)?; + self.put_call("", arguments, context)?; + } + crate::Expression::Call { + origin: crate::FunctionOrigin::External(ref name), + ref arguments, + } => match name.as_str() { + "atan2" | "cos" | "distance" | "length" | "mix" | "normalize" | "sin" => { + self.put_call(name, arguments, context)?; + } + "fclamp" => { + self.put_call("clamp", arguments, context)?; + } + other => return Err(Error::UnsupportedCall(other.to_owned())), + }, + crate::Expression::ArrayLength(expr) => match *self + .typifier + .get(expr, &context.module.types) + { + crate::TypeInner::Array { + size: crate::ArraySize::Constant(const_handle), + .. + } => { + self.put_constant(const_handle, context.module)?; + } + crate::TypeInner::Array { .. } => return Err(Error::UnsupportedDynamicArrayLength), + _ => return Err(Error::Validation), + }, + } + Ok(()) + } + + fn put_constant( + &mut self, + handle: Handle<crate::Constant>, + module: &crate::Module, + ) -> Result<(), Error> { + let constant = &module.constants[handle]; + match constant.inner { + crate::ConstantInner::Sint(value) => { + write!(self.out, "{}", value)?; + } + crate::ConstantInner::Uint(value) => { + write!(self.out, "{}", value)?; + } + crate::ConstantInner::Float(value) => { + write!(self.out, "{}", value)?; + if value.fract() == 0.0 { + write!(self.out, ".0")?; + } + } + crate::ConstantInner::Bool(value) => { + write!(self.out, "{}", value)?; + } + crate::ConstantInner::Composite(ref constituents) => { + let ty_name = &self.names[&NameKey::Type(constant.ty)]; + write!(self.out, "{}(", ty_name)?; + for (i, &handle) in constituents.iter().enumerate() { + if i != 0 { + write!(self.out, ", ")?; + } + self.put_constant(handle, module)?; + } + write!(self.out, ")")?; + } + } + Ok(()) + } + + fn put_block( + &mut self, + level: Level, + statements: &[crate::Statement], + context: &ExpressionContext, + return_value: Option<&str>, + ) -> Result<(), Error> { + for statement in statements { + log::trace!("statement[{}] {:?}", level.0, statement); + match *statement { + crate::Statement::Block(ref block) => { + if !block.is_empty() { + writeln!(self.out, "{}{{", level)?; + self.put_block(level.next(), block, context, return_value)?; + writeln!(self.out, "{}}}", level)?; + } + } + crate::Statement::If { + condition, + ref accept, + ref reject, + } => { + write!(self.out, "{}if (", level)?; + self.put_expression(condition, context)?; + writeln!(self.out, ") {{")?; + self.put_block(level.next(), accept, context, return_value)?; + if !reject.is_empty() { + writeln!(self.out, "{}}} else {{", level)?; + self.put_block(level.next(), reject, context, return_value)?; + } + writeln!(self.out, "{}}}", level)?; + } + crate::Statement::Switch { + selector, + ref cases, + ref default, + } => { + write!(self.out, "{}switch(", level)?; + self.put_expression(selector, context)?; + writeln!(self.out, ") {{")?; + let lcase = level.next(); + for (&value, &(ref block, ref fall_through)) in cases.iter() { + writeln!(self.out, "{}case {}: {{", lcase, value)?; + self.put_block(lcase.next(), block, context, return_value)?; + if fall_through.is_none() { + writeln!(self.out, "{}break;", lcase.next())?; + } + writeln!(self.out, "{}}}", lcase)?; + } + writeln!(self.out, "{}default: {{", lcase)?; + self.put_block(lcase.next(), default, context, return_value)?; + writeln!(self.out, "{}}}", lcase)?; + writeln!(self.out, "{}}}", level)?; + } + crate::Statement::Loop { + ref body, + ref continuing, + } => { + writeln!(self.out, "{}while(true) {{", level)?; + self.put_block(level.next(), body, context, return_value)?; + if !continuing.is_empty() { + //TODO + } + writeln!(self.out, "{}}}", level)?; + } + crate::Statement::Break => { + writeln!(self.out, "{}break;", level)?; + } + crate::Statement::Continue => { + writeln!(self.out, "{}continue;", level)?; + } + crate::Statement::Return { + value: Some(expr_handle), + } => { + write!(self.out, "{}return ", level)?; + self.put_expression(expr_handle, context)?; + writeln!(self.out, ";")?; + } + crate::Statement::Return { value: None } => { + if let Some(string) = return_value { + writeln!(self.out, "{}return {};", level, string)?; + } + } + crate::Statement::Kill => { + writeln!(self.out, "{}discard_fragment();", level)?; + } + crate::Statement::Store { pointer, value } => { + //write!(self.out, "\t*")?; + write!(self.out, "{}", level)?; + self.put_expression(pointer, context)?; + write!(self.out, " = ")?; + self.put_expression(value, context)?; + writeln!(self.out, ";")?; + } + } + } + Ok(()) + } + + pub fn write(&mut self, module: &crate::Module, options: &Options) -> Result<(), Error> { + self.names.clear(); + Namer::process(module, RESERVED, &mut self.names); + + writeln!(self.out, "#include <metal_stdlib>")?; + writeln!(self.out, "#include <simd/simd.h>")?; + + writeln!(self.out)?; + self.write_type_defs(module)?; + + writeln!(self.out)?; + self.write_functions(module, options)?; + + Ok(()) + } + + fn write_type_defs(&mut self, module: &crate::Module) -> Result<(), Error> { + for (handle, ty) in module.types.iter() { + let name = &self.names[&NameKey::Type(handle)]; + match ty.inner { + crate::TypeInner::Scalar { kind, .. } => { + write!(self.out, "typedef {} {}", scalar_kind_string(kind), name)?; + } + crate::TypeInner::Vector { size, kind, .. } => { + write!( + self.out, + "typedef {}{} {}", + scalar_kind_string(kind), + vector_size_string(size), + name + )?; + } + crate::TypeInner::Matrix { columns, rows, .. } => { + write!( + self.out, + "typedef {}{}x{} {}", + scalar_kind_string(crate::ScalarKind::Float), + vector_size_string(columns), + vector_size_string(rows), + name + )?; + } + crate::TypeInner::Pointer { base, class } => { + use crate::StorageClass as Sc; + let base_name = &self.names[&NameKey::Type(base)]; + let class_name = match class { + Sc::Input | Sc::Output => continue, + Sc::Uniform => "constant", + Sc::Storage => "device", + Sc::Handle + | Sc::Private + | Sc::Function + | Sc::WorkGroup + | Sc::PushConstant => "", + }; + write!(self.out, "typedef {} {} *{}", class_name, base_name, name)?; + } + crate::TypeInner::Array { + base, + size, + stride: _, + } => { + let base_name = &self.names[&NameKey::Type(base)]; + write!(self.out, "typedef {} {}[", base_name, name)?; + match size { + crate::ArraySize::Constant(const_handle) => { + self.put_constant(const_handle, module)?; + write!(self.out, "]")?; + } + crate::ArraySize::Dynamic => write!(self.out, "1]")?, + } + } + crate::TypeInner::Struct { ref members } => { + writeln!(self.out, "struct {} {{", name)?; + for (index, member) in members.iter().enumerate() { + let member_name = &self.names[&NameKey::StructMember(handle, index as u32)]; + let base_name = &self.names[&NameKey::Type(member.ty)]; + write!(self.out, "\t{} {}", base_name, member_name)?; + match member.origin { + crate::MemberOrigin::Empty => {} + crate::MemberOrigin::BuiltIn(built_in) => { + ResolvedBinding::BuiltIn(built_in) + .try_fmt_decorated(&mut self.out, "")?; + } + crate::MemberOrigin::Offset(_) => { + //TODO + } + } + writeln!(self.out, ";")?; + } + write!(self.out, "}}")?; + } + crate::TypeInner::Image { + dim, + arrayed, + class, + } => { + let dim_str = match dim { + crate::ImageDimension::D1 => "1d", + crate::ImageDimension::D2 => "2d", + crate::ImageDimension::D3 => "3d", + crate::ImageDimension::Cube => "Cube", + }; + let (texture_str, msaa_str, kind, access) = match class { + crate::ImageClass::Sampled { kind, multi } => { + ("texture", if multi { "_ms" } else { "" }, kind, "sample") + } + crate::ImageClass::Depth => { + ("depth", "", crate::ScalarKind::Float, "sample") + } + crate::ImageClass::Storage(format) => { + let (_, global) = module + .global_variables + .iter() + .find(|(_, var)| var.ty == handle) + .expect("Unable to find a global variable using the image type"); + let access = if global + .storage_access + .contains(crate::StorageAccess::LOAD | crate::StorageAccess::STORE) + { + "read_write" + } else if global.storage_access.contains(crate::StorageAccess::STORE) { + "write" + } else if global.storage_access.contains(crate::StorageAccess::LOAD) { + "read" + } else { + return Err(Error::InvalidImageAccess(global.storage_access)); + }; + ("texture", "", format.into(), access) + } + }; + let base_name = scalar_kind_string(kind); + let array_str = if arrayed { "_array" } else { "" }; + write!( + self.out, + "typedef {}{}{}{}<{}, access::{}> {}", + texture_str, dim_str, msaa_str, array_str, base_name, access, name + )?; + } + crate::TypeInner::Sampler { comparison: _ } => { + write!(self.out, "typedef sampler {}", name)?; + } + } + writeln!(self.out, ";")?; + } + Ok(()) + } + + fn write_functions(&mut self, module: &crate::Module, options: &Options) -> Result<(), Error> { + for (fun_handle, fun) in module.functions.iter() { + self.typifier.resolve_all( + &fun.expressions, + &module.types, + &ResolveContext { + constants: &module.constants, + global_vars: &module.global_variables, + local_vars: &fun.local_variables, + functions: &module.functions, + arguments: &fun.arguments, + }, + )?; + + let fun_name = &self.names[&NameKey::Function(fun_handle)]; + let result_type_name = match fun.return_type { + Some(ret_ty) => &self.names[&NameKey::Type(ret_ty)], + None => "void", + }; + writeln!(self.out, "{} {}(", result_type_name, fun_name)?; + + for (index, arg) in fun.arguments.iter().enumerate() { + let name = &self.names[&NameKey::FunctionArgument(fun_handle, index as u32)]; + let param_type_name = &self.names[&NameKey::Type(arg.ty)]; + let separator = separate(index + 1 == fun.arguments.len()); + writeln!(self.out, "\t{} {}{}", param_type_name, name, separator)?; + } + writeln!(self.out, ") {{")?; + + for (local_handle, local) in fun.local_variables.iter() { + let ty_name = &self.names[&NameKey::Type(local.ty)]; + let local_name = &self.names[&NameKey::FunctionLocal(fun_handle, local_handle)]; + write!(self.out, "\t{} {}", ty_name, local_name)?; + if let Some(value) = local.init { + write!(self.out, " = ")?; + self.put_constant(value, module)?; + } + writeln!(self.out, ";")?; + } + + let context = ExpressionContext { + function: fun, + origin: FunctionOrigin::Handle(fun_handle), + module, + }; + self.put_block(Level(1), &fun.body, &context, None)?; + writeln!(self.out, "}}")?; + } + + for (ep_index, (&(stage, _), ep)) in module.entry_points.iter().enumerate() { + let fun = &ep.function; + self.typifier.resolve_all( + &fun.expressions, + &module.types, + &ResolveContext { + constants: &module.constants, + global_vars: &module.global_variables, + local_vars: &fun.local_variables, + functions: &module.functions, + arguments: &fun.arguments, + }, + )?; + + // find the entry point(s) and inputs/outputs + let mut last_used_global = None; + for ((handle, var), &usage) in module.global_variables.iter().zip(&fun.global_usage) { + match var.class { + crate::StorageClass::Input => { + if let Some(crate::Binding::Location(_)) = var.binding { + continue; + } + } + crate::StorageClass::Output => continue, + _ => {} + } + if !usage.is_empty() { + last_used_global = Some(handle); + } + } + + let fun_name = &self.names[&NameKey::EntryPoint(ep_index as _)]; + let output_name = format!("{}Output", fun_name); + let location_input_name = format!("{}Input", fun_name); + + let (em_str, in_mode, out_mode) = match stage { + crate::ShaderStage::Vertex => ( + "vertex", + LocationMode::VertexInput, + LocationMode::Intermediate, + ), + crate::ShaderStage::Fragment { .. } => ( + "fragment", + LocationMode::Intermediate, + LocationMode::FragmentOutput, + ), + crate::ShaderStage::Compute { .. } => { + ("kernel", LocationMode::Uniform, LocationMode::Uniform) + } + }; + + let return_value = match stage { + crate::ShaderStage::Vertex | crate::ShaderStage::Fragment => { + // make dedicated input/output structs + writeln!(self.out, "struct {} {{", location_input_name)?; + + for ((handle, var), &usage) in + module.global_variables.iter().zip(&fun.global_usage) + { + if var.class != crate::StorageClass::Input + || !usage.contains(crate::GlobalUse::LOAD) + { + continue; + } + // if it's a struct, lift all the built-in contents up to the root + if let crate::TypeInner::Struct { ref members } = module.types[var.ty].inner + { + for (index, member) in members.iter().enumerate() { + if let crate::MemberOrigin::BuiltIn(built_in) = member.origin { + let name = + &self.names[&NameKey::StructMember(var.ty, index as u32)]; + let ty_name = &self.names[&NameKey::Type(member.ty)]; + write!(self.out, "\t{} {}", ty_name, name)?; + ResolvedBinding::BuiltIn(built_in) + .try_fmt_decorated(&mut self.out, ";\n")?; + } + } + } else if let Some(ref binding @ crate::Binding::Location(_)) = var.binding + { + let tyvar = TypedGlobalVariable { + module, + names: &self.names, + handle, + usage: crate::GlobalUse::empty(), + }; + let resolved = options.resolve_binding(stage, binding, in_mode)?; + + write!(self.out, "\t")?; + tyvar.try_fmt(&mut self.out)?; + resolved.try_fmt_decorated(&mut self.out, ";\n")?; + } + } + writeln!(self.out, "}};")?; + + writeln!(self.out, "struct {} {{", output_name)?; + for ((handle, var), &usage) in + module.global_variables.iter().zip(&fun.global_usage) + { + if var.class != crate::StorageClass::Output + || !usage.contains(crate::GlobalUse::STORE) + { + continue; + } + // if it's a struct, lift all the built-in contents up to the root + if let crate::TypeInner::Struct { ref members } = module.types[var.ty].inner + { + for (index, member) in members.iter().enumerate() { + let name = + &self.names[&NameKey::StructMember(var.ty, index as u32)]; + let ty_name = &self.names[&NameKey::Type(member.ty)]; + match member.origin { + crate::MemberOrigin::Empty => {} + crate::MemberOrigin::BuiltIn(built_in) => { + write!(self.out, "\t{} {}", ty_name, name)?; + ResolvedBinding::BuiltIn(built_in) + .try_fmt_decorated(&mut self.out, ";\n")?; + } + crate::MemberOrigin::Offset(_) => { + //TODO + } + } + } + } else { + let tyvar = TypedGlobalVariable { + module, + names: &self.names, + handle, + usage: crate::GlobalUse::empty(), + }; + write!(self.out, "\t")?; + tyvar.try_fmt(&mut self.out)?; + if let Some(ref binding) = var.binding { + let resolved = options.resolve_binding(stage, binding, out_mode)?; + resolved.try_fmt_decorated(&mut self.out, "")?; + } + writeln!(self.out, ";")?; + } + } + writeln!(self.out, "}};")?; + + writeln!(self.out, "{} {} {}(", em_str, output_name, fun_name)?; + let separator = separate(last_used_global.is_none()); + writeln!( + self.out, + "\t{} {} [[stage_in]]{}", + location_input_name, LOCATION_INPUT_STRUCT_NAME, separator + )?; + + Some(OUTPUT_STRUCT_NAME) + } + crate::ShaderStage::Compute => { + writeln!(self.out, "{} void {}(", em_str, fun_name)?; + None + } + }; + + for ((handle, var), &usage) in module.global_variables.iter().zip(&fun.global_usage) { + if usage.is_empty() || var.class == crate::StorageClass::Output { + continue; + } + if var.class == crate::StorageClass::Input { + if let Some(crate::Binding::Location(_)) = var.binding { + // location inputs are put into a separate struct + continue; + } + } + let loc_mode = match (stage, var.class) { + (crate::ShaderStage::Vertex, crate::StorageClass::Input) => { + LocationMode::VertexInput + } + (crate::ShaderStage::Vertex, crate::StorageClass::Output) + | (crate::ShaderStage::Fragment { .. }, crate::StorageClass::Input) => { + LocationMode::Intermediate + } + (crate::ShaderStage::Fragment { .. }, crate::StorageClass::Output) => { + LocationMode::FragmentOutput + } + _ => LocationMode::Uniform, + }; + let resolved = + options.resolve_binding(stage, var.binding.as_ref().unwrap(), loc_mode)?; + let tyvar = TypedGlobalVariable { + module, + names: &self.names, + handle, + usage, + }; + let separator = separate(last_used_global == Some(handle)); + write!(self.out, "\t")?; + tyvar.try_fmt(&mut self.out)?; + resolved.try_fmt_decorated(&mut self.out, separator)?; + if let Some(value) = var.init { + write!(self.out, " = ")?; + self.put_constant(value, module)?; + } + writeln!(self.out)?; + } + writeln!(self.out, ") {{")?; + + match stage { + crate::ShaderStage::Vertex | crate::ShaderStage::Fragment => { + writeln!(self.out, "\t{} {};", output_name, OUTPUT_STRUCT_NAME)?; + } + crate::ShaderStage::Compute => {} + } + for (local_handle, local) in fun.local_variables.iter() { + let name = &self.names[&NameKey::EntryPointLocal(ep_index as _, local_handle)]; + let ty_name = &self.names[&NameKey::Type(local.ty)]; + write!(self.out, "\t{} {}", ty_name, name)?; + if let Some(value) = local.init { + write!(self.out, " = ")?; + self.put_constant(value, module)?; + } + writeln!(self.out, ";")?; + } + + let context = ExpressionContext { + function: fun, + origin: FunctionOrigin::EntryPoint(ep_index as _), + module, + }; + self.put_block(Level(1), &fun.body, &context, return_value)?; + writeln!(self.out, "}}")?; + } + + Ok(()) + } +} diff --git a/third_party/rust/naga/src/back/spv/helpers.rs b/third_party/rust/naga/src/back/spv/helpers.rs new file mode 100644 index 0000000000..5facbe8b69 --- /dev/null +++ b/third_party/rust/naga/src/back/spv/helpers.rs @@ -0,0 +1,20 @@ +use spirv::Word; + +pub(crate) fn bytes_to_words(bytes: &[u8]) -> Vec<Word> { + bytes + .chunks(4) + .map(|chars| chars.iter().rev().fold(0u32, |u, c| (u << 8) | *c as u32)) + .collect() +} + +pub(crate) fn string_to_words(input: &str) -> Vec<Word> { + let bytes = input.as_bytes(); + let mut words = bytes_to_words(bytes); + + if bytes.len() % 4 == 0 { + // nul-termination + words.push(0x0u32); + } + + words +} diff --git a/third_party/rust/naga/src/back/spv/instructions.rs b/third_party/rust/naga/src/back/spv/instructions.rs new file mode 100644 index 0000000000..ab8e56844a --- /dev/null +++ b/third_party/rust/naga/src/back/spv/instructions.rs @@ -0,0 +1,708 @@ +use crate::back::spv::{helpers, Instruction}; +use spirv::{Op, Word}; + +pub(super) enum Signedness { + Unsigned = 0, + Signed = 1, +} + +// +// Debug Instructions +// + +pub(super) fn instruction_source( + source_language: spirv::SourceLanguage, + version: u32, +) -> Instruction { + let mut instruction = Instruction::new(Op::Source); + instruction.add_operand(source_language as u32); + instruction.add_operands(helpers::bytes_to_words(&version.to_le_bytes())); + instruction +} + +pub(super) fn instruction_name(target_id: Word, name: &str) -> Instruction { + let mut instruction = Instruction::new(Op::Name); + instruction.add_operand(target_id); + instruction.add_operands(helpers::string_to_words(name)); + instruction +} + +// +// Annotation Instructions +// + +pub(super) fn instruction_decorate( + target_id: Word, + decoration: spirv::Decoration, + operands: &[Word], +) -> Instruction { + let mut instruction = Instruction::new(Op::Decorate); + instruction.add_operand(target_id); + instruction.add_operand(decoration as u32); + + for operand in operands { + instruction.add_operand(*operand) + } + + instruction +} + +// +// Extension Instructions +// + +pub(super) fn instruction_ext_inst_import(id: Word, name: &str) -> Instruction { + let mut instruction = Instruction::new(Op::ExtInstImport); + instruction.set_result(id); + instruction.add_operands(helpers::string_to_words(name)); + instruction +} + +// +// Mode-Setting Instructions +// + +pub(super) fn instruction_memory_model( + addressing_model: spirv::AddressingModel, + memory_model: spirv::MemoryModel, +) -> Instruction { + let mut instruction = Instruction::new(Op::MemoryModel); + instruction.add_operand(addressing_model as u32); + instruction.add_operand(memory_model as u32); + instruction +} + +pub(super) fn instruction_entry_point( + execution_model: spirv::ExecutionModel, + entry_point_id: Word, + name: &str, + interface_ids: &[Word], +) -> Instruction { + let mut instruction = Instruction::new(Op::EntryPoint); + instruction.add_operand(execution_model as u32); + instruction.add_operand(entry_point_id); + instruction.add_operands(helpers::string_to_words(name)); + + for interface_id in interface_ids { + instruction.add_operand(*interface_id); + } + + instruction +} + +pub(super) fn instruction_execution_mode( + entry_point_id: Word, + execution_mode: spirv::ExecutionMode, +) -> Instruction { + let mut instruction = Instruction::new(Op::ExecutionMode); + instruction.add_operand(entry_point_id); + instruction.add_operand(execution_mode as u32); + instruction +} + +pub(super) fn instruction_capability(capability: spirv::Capability) -> Instruction { + let mut instruction = Instruction::new(Op::Capability); + instruction.add_operand(capability as u32); + instruction +} + +// +// Type-Declaration Instructions +// + +pub(super) fn instruction_type_void(id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::TypeVoid); + instruction.set_result(id); + instruction +} + +pub(super) fn instruction_type_bool(id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::TypeBool); + instruction.set_result(id); + instruction +} + +pub(super) fn instruction_type_int(id: Word, width: Word, signedness: Signedness) -> Instruction { + let mut instruction = Instruction::new(Op::TypeInt); + instruction.set_result(id); + instruction.add_operand(width); + instruction.add_operand(signedness as u32); + instruction +} + +pub(super) fn instruction_type_float(id: Word, width: Word) -> Instruction { + let mut instruction = Instruction::new(Op::TypeFloat); + instruction.set_result(id); + instruction.add_operand(width); + instruction +} + +pub(super) fn instruction_type_vector( + id: Word, + component_type_id: Word, + component_count: crate::VectorSize, +) -> Instruction { + let mut instruction = Instruction::new(Op::TypeVector); + instruction.set_result(id); + instruction.add_operand(component_type_id); + instruction.add_operand(component_count as u32); + instruction +} + +pub(super) fn instruction_type_matrix( + id: Word, + column_type_id: Word, + column_count: crate::VectorSize, +) -> Instruction { + let mut instruction = Instruction::new(Op::TypeMatrix); + instruction.set_result(id); + instruction.add_operand(column_type_id); + instruction.add_operand(column_count as u32); + instruction +} + +pub(super) fn instruction_type_image( + id: Word, + sampled_type_id: Word, + dim: spirv::Dim, + arrayed: bool, + image_class: crate::ImageClass, +) -> Instruction { + let mut instruction = Instruction::new(Op::TypeImage); + instruction.set_result(id); + instruction.add_operand(sampled_type_id); + instruction.add_operand(dim as u32); + + instruction.add_operand(match image_class { + crate::ImageClass::Depth => 1, + _ => 0, + }); + instruction.add_operand(arrayed as u32); + instruction.add_operand(match image_class { + crate::ImageClass::Sampled { multi: true, .. } => 1, + _ => 0, + }); + instruction.add_operand(match image_class { + crate::ImageClass::Sampled { .. } => 1, + _ => 0, + }); + + let format = match image_class { + crate::ImageClass::Storage(format) => match format { + crate::StorageFormat::R8Unorm => spirv::ImageFormat::R8, + crate::StorageFormat::R8Snorm => spirv::ImageFormat::R8Snorm, + crate::StorageFormat::R8Uint => spirv::ImageFormat::R8ui, + crate::StorageFormat::R8Sint => spirv::ImageFormat::R8i, + crate::StorageFormat::R16Uint => spirv::ImageFormat::R16ui, + crate::StorageFormat::R16Sint => spirv::ImageFormat::R16i, + crate::StorageFormat::R16Float => spirv::ImageFormat::R16f, + crate::StorageFormat::Rg8Unorm => spirv::ImageFormat::Rg8, + crate::StorageFormat::Rg8Snorm => spirv::ImageFormat::Rg8Snorm, + crate::StorageFormat::Rg8Uint => spirv::ImageFormat::Rg8ui, + crate::StorageFormat::Rg8Sint => spirv::ImageFormat::Rg8i, + crate::StorageFormat::R32Uint => spirv::ImageFormat::R32ui, + crate::StorageFormat::R32Sint => spirv::ImageFormat::R32i, + crate::StorageFormat::R32Float => spirv::ImageFormat::R32f, + crate::StorageFormat::Rg16Uint => spirv::ImageFormat::Rg16ui, + crate::StorageFormat::Rg16Sint => spirv::ImageFormat::Rg16i, + crate::StorageFormat::Rg16Float => spirv::ImageFormat::Rg16f, + crate::StorageFormat::Rgba8Unorm => spirv::ImageFormat::Rgba8, + crate::StorageFormat::Rgba8Snorm => spirv::ImageFormat::Rgba8Snorm, + crate::StorageFormat::Rgba8Uint => spirv::ImageFormat::Rgba8ui, + crate::StorageFormat::Rgba8Sint => spirv::ImageFormat::Rgba8i, + crate::StorageFormat::Rgb10a2Unorm => spirv::ImageFormat::Rgb10a2ui, + crate::StorageFormat::Rg11b10Float => spirv::ImageFormat::R11fG11fB10f, + crate::StorageFormat::Rg32Uint => spirv::ImageFormat::Rg32ui, + crate::StorageFormat::Rg32Sint => spirv::ImageFormat::Rg32i, + crate::StorageFormat::Rg32Float => spirv::ImageFormat::Rg32f, + crate::StorageFormat::Rgba16Uint => spirv::ImageFormat::Rgba16ui, + crate::StorageFormat::Rgba16Sint => spirv::ImageFormat::Rgba16i, + crate::StorageFormat::Rgba16Float => spirv::ImageFormat::Rgba16f, + crate::StorageFormat::Rgba32Uint => spirv::ImageFormat::Rgba32ui, + crate::StorageFormat::Rgba32Sint => spirv::ImageFormat::Rgba32i, + crate::StorageFormat::Rgba32Float => spirv::ImageFormat::Rgba32f, + }, + _ => spirv::ImageFormat::Unknown, + }; + + instruction.add_operand(format as u32); + instruction +} + +pub(super) fn instruction_type_sampler(id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::TypeSampler); + instruction.set_result(id); + instruction +} + +pub(super) fn instruction_type_sampled_image(id: Word, image_type_id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::TypeSampledImage); + instruction.set_result(id); + instruction.add_operand(image_type_id); + instruction +} + +pub(super) fn instruction_type_array( + id: Word, + element_type_id: Word, + length_id: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::TypeArray); + instruction.set_result(id); + instruction.add_operand(element_type_id); + instruction.add_operand(length_id); + instruction +} + +pub(super) fn instruction_type_runtime_array(id: Word, element_type_id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::TypeRuntimeArray); + instruction.set_result(id); + instruction.add_operand(element_type_id); + instruction +} + +pub(super) fn instruction_type_struct(id: Word, member_ids: &[Word]) -> Instruction { + let mut instruction = Instruction::new(Op::TypeStruct); + instruction.set_result(id); + + for member_id in member_ids { + instruction.add_operand(*member_id) + } + + instruction +} + +pub(super) fn instruction_type_pointer( + id: Word, + storage_class: spirv::StorageClass, + type_id: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::TypePointer); + instruction.set_result(id); + instruction.add_operand(storage_class as u32); + instruction.add_operand(type_id); + instruction +} + +pub(super) fn instruction_type_function( + id: Word, + return_type_id: Word, + parameter_ids: &[Word], +) -> Instruction { + let mut instruction = Instruction::new(Op::TypeFunction); + instruction.set_result(id); + instruction.add_operand(return_type_id); + + for parameter_id in parameter_ids { + instruction.add_operand(*parameter_id); + } + + instruction +} + +// +// Constant-Creation Instructions +// + +pub(super) fn instruction_constant_true(result_type_id: Word, id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::ConstantTrue); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction +} + +pub(super) fn instruction_constant_false(result_type_id: Word, id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::ConstantFalse); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction +} + +pub(super) fn instruction_constant(result_type_id: Word, id: Word, values: &[Word]) -> Instruction { + let mut instruction = Instruction::new(Op::Constant); + instruction.set_type(result_type_id); + instruction.set_result(id); + + for value in values { + instruction.add_operand(*value); + } + + instruction +} + +pub(super) fn instruction_constant_composite( + result_type_id: Word, + id: Word, + constituent_ids: &[Word], +) -> Instruction { + let mut instruction = Instruction::new(Op::ConstantComposite); + instruction.set_type(result_type_id); + instruction.set_result(id); + + for constituent_id in constituent_ids { + instruction.add_operand(*constituent_id); + } + + instruction +} + +// +// Memory Instructions +// + +pub(super) fn instruction_variable( + result_type_id: Word, + id: Word, + storage_class: spirv::StorageClass, + initializer_id: Option<Word>, +) -> Instruction { + let mut instruction = Instruction::new(Op::Variable); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(storage_class as u32); + + if let Some(initializer_id) = initializer_id { + instruction.add_operand(initializer_id); + } + + instruction +} + +pub(super) fn instruction_load( + result_type_id: Word, + id: Word, + pointer_type_id: Word, + memory_access: Option<spirv::MemoryAccess>, +) -> Instruction { + let mut instruction = Instruction::new(Op::Load); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(pointer_type_id); + + if let Some(memory_access) = memory_access { + instruction.add_operand(memory_access.bits()); + } + + instruction +} + +pub(super) fn instruction_store( + pointer_type_id: Word, + object_id: Word, + memory_access: Option<spirv::MemoryAccess>, +) -> Instruction { + let mut instruction = Instruction::new(Op::Store); + instruction.add_operand(pointer_type_id); + instruction.add_operand(object_id); + + if let Some(memory_access) = memory_access { + instruction.add_operand(memory_access.bits()); + } + + instruction +} + +pub(super) fn instruction_access_chain( + result_type_id: Word, + id: Word, + base_id: Word, + index_ids: &[Word], +) -> Instruction { + let mut instruction = Instruction::new(Op::AccessChain); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(base_id); + + for index_id in index_ids { + instruction.add_operand(*index_id); + } + + instruction +} + +// +// Function Instructions +// + +pub(super) fn instruction_function( + return_type_id: Word, + id: Word, + function_control: spirv::FunctionControl, + function_type_id: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::Function); + instruction.set_type(return_type_id); + instruction.set_result(id); + instruction.add_operand(function_control.bits()); + instruction.add_operand(function_type_id); + instruction +} + +pub(super) fn instruction_function_parameter(result_type_id: Word, id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::FunctionParameter); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction +} + +pub(super) fn instruction_function_end() -> Instruction { + Instruction::new(Op::FunctionEnd) +} + +pub(super) fn instruction_function_call( + result_type_id: Word, + id: Word, + function_id: Word, + argument_ids: &[Word], +) -> Instruction { + let mut instruction = Instruction::new(Op::FunctionCall); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(function_id); + + for argument_id in argument_ids { + instruction.add_operand(*argument_id); + } + + instruction +} + +// +// Image Instructions +// +pub(super) fn instruction_sampled_image( + result_type_id: Word, + id: Word, + image: Word, + sampler: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::SampledImage); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(image); + instruction.add_operand(sampler); + instruction +} + +pub(super) fn instruction_image_sample_implicit_lod( + result_type_id: Word, + id: Word, + sampled_image: Word, + coordinates: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::ImageSampleImplicitLod); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(sampled_image); + instruction.add_operand(coordinates); + instruction +} + +// +// Conversion Instructions +// +pub(super) fn instruction_unary( + op: Op, + result_type_id: Word, + id: Word, + value: Word, +) -> Instruction { + let mut instruction = Instruction::new(op); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(value); + instruction +} + +// +// Composite Instructions +// + +pub(super) fn instruction_composite_construct( + result_type_id: Word, + id: Word, + constituent_ids: &[Word], +) -> Instruction { + let mut instruction = Instruction::new(Op::CompositeConstruct); + instruction.set_type(result_type_id); + instruction.set_result(id); + + for constituent_id in constituent_ids { + instruction.add_operand(*constituent_id); + } + + instruction +} + +// +// Arithmetic Instructions +// +fn instruction_binary( + op: Op, + result_type_id: Word, + id: Word, + operand_1: Word, + operand_2: Word, +) -> Instruction { + let mut instruction = Instruction::new(op); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(operand_1); + instruction.add_operand(operand_2); + instruction +} + +pub(super) fn instruction_i_sub( + result_type_id: Word, + id: Word, + operand_1: Word, + operand_2: Word, +) -> Instruction { + instruction_binary(Op::ISub, result_type_id, id, operand_1, operand_2) +} + +pub(super) fn instruction_f_sub( + result_type_id: Word, + id: Word, + operand_1: Word, + operand_2: Word, +) -> Instruction { + instruction_binary(Op::FSub, result_type_id, id, operand_1, operand_2) +} + +pub(super) fn instruction_i_mul( + result_type_id: Word, + id: Word, + operand_1: Word, + operand_2: Word, +) -> Instruction { + instruction_binary(Op::IMul, result_type_id, id, operand_1, operand_2) +} + +pub(super) fn instruction_f_mul( + result_type_id: Word, + id: Word, + operand_1: Word, + operand_2: Word, +) -> Instruction { + instruction_binary(Op::FMul, result_type_id, id, operand_1, operand_2) +} + +pub(super) fn instruction_vector_times_scalar( + result_type_id: Word, + id: Word, + vector_type_id: Word, + scalar_type_id: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::VectorTimesScalar); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(vector_type_id); + instruction.add_operand(scalar_type_id); + instruction +} + +pub(super) fn instruction_matrix_times_scalar( + result_type_id: Word, + id: Word, + matrix_id: Word, + scalar_id: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::MatrixTimesScalar); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(matrix_id); + instruction.add_operand(scalar_id); + instruction +} + +pub(super) fn instruction_vector_times_matrix( + result_type_id: Word, + id: Word, + vector_id: Word, + matrix_id: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::VectorTimesMatrix); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(vector_id); + instruction.add_operand(matrix_id); + instruction +} + +pub(super) fn instruction_matrix_times_vector( + result_type_id: Word, + id: Word, + matrix_id: Word, + vector_id: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::MatrixTimesVector); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(matrix_id); + instruction.add_operand(vector_id); + instruction +} + +pub(super) fn instruction_matrix_times_matrix( + result_type_id: Word, + id: Word, + left_matrix: Word, + right_matrix: Word, +) -> Instruction { + let mut instruction = Instruction::new(Op::MatrixTimesMatrix); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(left_matrix); + instruction.add_operand(right_matrix); + instruction +} + +// +// Bit Instructions +// + +pub(super) fn instruction_bitwise_and( + result_type_id: Word, + id: Word, + operand_1: Word, + operand_2: Word, +) -> Instruction { + instruction_binary(Op::BitwiseAnd, result_type_id, id, operand_1, operand_2) +} + +// +// Relational and Logical Instructions +// + +// +// Derivative Instructions +// + +// +// Control-Flow Instructions +// + +pub(super) fn instruction_label(id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::Label); + instruction.set_result(id); + instruction +} + +pub(super) fn instruction_return() -> Instruction { + Instruction::new(Op::Return) +} + +pub(super) fn instruction_return_value(value_id: Word) -> Instruction { + let mut instruction = Instruction::new(Op::ReturnValue); + instruction.add_operand(value_id); + instruction +} + +// +// Atomic Instructions +// + +// +// Primitive Instructions +// diff --git a/third_party/rust/naga/src/back/spv/layout.rs b/third_party/rust/naga/src/back/spv/layout.rs new file mode 100644 index 0000000000..006e785317 --- /dev/null +++ b/third_party/rust/naga/src/back/spv/layout.rs @@ -0,0 +1,91 @@ +use crate::back::spv::{Instruction, LogicalLayout, PhysicalLayout}; +use spirv::*; +use std::iter; + +impl PhysicalLayout { + pub(super) fn new(header: &crate::Header) -> Self { + let version: Word = ((header.version.0 as u32) << 16) + | ((header.version.1 as u32) << 8) + | header.version.2 as u32; + + PhysicalLayout { + magic_number: MAGIC_NUMBER, + version, + generator: header.generator, + bound: 0, + instruction_schema: 0x0u32, + } + } + + pub(super) fn in_words(&self, sink: &mut impl Extend<Word>) { + sink.extend(iter::once(self.magic_number)); + sink.extend(iter::once(self.version)); + sink.extend(iter::once(self.generator)); + sink.extend(iter::once(self.bound)); + sink.extend(iter::once(self.instruction_schema)); + } + + pub(super) fn supports_storage_buffers(&self) -> bool { + self.version >= 0x10300 + } +} + +impl LogicalLayout { + pub(super) fn in_words(&self, sink: &mut impl Extend<Word>) { + sink.extend(self.capabilities.iter().cloned()); + sink.extend(self.extensions.iter().cloned()); + sink.extend(self.ext_inst_imports.iter().cloned()); + sink.extend(self.memory_model.iter().cloned()); + sink.extend(self.entry_points.iter().cloned()); + sink.extend(self.execution_modes.iter().cloned()); + sink.extend(self.debugs.iter().cloned()); + sink.extend(self.annotations.iter().cloned()); + sink.extend(self.declarations.iter().cloned()); + sink.extend(self.function_declarations.iter().cloned()); + sink.extend(self.function_definitions.iter().cloned()); + } +} + +impl Instruction { + pub(super) fn new(op: Op) -> Self { + Instruction { + op, + wc: 1, // Always start at 1 for the first word (OP + WC), + type_id: None, + result_id: None, + operands: vec![], + } + } + + #[allow(clippy::panic)] + pub(super) fn set_type(&mut self, id: Word) { + assert!(self.type_id.is_none(), "Type can only be set once"); + self.type_id = Some(id); + self.wc += 1; + } + + #[allow(clippy::panic)] + pub(super) fn set_result(&mut self, id: Word) { + assert!(self.result_id.is_none(), "Result can only be set once"); + self.result_id = Some(id); + self.wc += 1; + } + + pub(super) fn add_operand(&mut self, operand: Word) { + self.operands.push(operand); + self.wc += 1; + } + + pub(super) fn add_operands(&mut self, operands: Vec<Word>) { + for operand in operands.into_iter() { + self.add_operand(operand) + } + } + + pub(super) fn to_words(&self, sink: &mut impl Extend<Word>) { + sink.extend(Some((self.wc << 16 | self.op as u32) as u32)); + sink.extend(self.type_id); + sink.extend(self.result_id); + sink.extend(self.operands.iter().cloned()); + } +} diff --git a/third_party/rust/naga/src/back/spv/layout_tests.rs b/third_party/rust/naga/src/back/spv/layout_tests.rs new file mode 100644 index 0000000000..37024b238f --- /dev/null +++ b/third_party/rust/naga/src/back/spv/layout_tests.rs @@ -0,0 +1,166 @@ +use crate::back::spv::test_framework::*; +use crate::back::spv::{helpers, Instruction, LogicalLayout, PhysicalLayout}; +use crate::Header; +use spirv::*; + +#[test] +fn test_physical_layout_in_words() { + let header = Header { + generator: 0, + version: (1, 2, 3), + }; + let bound = 5; + + let mut output = vec![]; + let mut layout = PhysicalLayout::new(&header); + layout.bound = bound; + + layout.in_words(&mut output); + + assert_eq!(output[0], spirv::MAGIC_NUMBER); + assert_eq!( + output[1], + to_word(&[header.version.0, header.version.1, header.version.2, 1]) + ); + assert_eq!(output[2], 0); + assert_eq!(output[3], bound); + assert_eq!(output[4], 0); +} + +#[test] +fn test_logical_layout_in_words() { + let mut output = vec![]; + let mut layout = LogicalLayout::default(); + let layout_vectors = 11; + let mut instructions = Vec::with_capacity(layout_vectors); + + let vector_names = &[ + "Capabilities", + "Extensions", + "External Instruction Imports", + "Memory Model", + "Entry Points", + "Execution Modes", + "Debugs", + "Annotations", + "Declarations", + "Function Declarations", + "Function Definitions", + ]; + + for i in 0..layout_vectors { + let mut dummy_instruction = Instruction::new(Op::Constant); + dummy_instruction.set_type((i + 1) as u32); + dummy_instruction.set_result((i + 2) as u32); + dummy_instruction.add_operand((i + 3) as u32); + dummy_instruction.add_operands(helpers::string_to_words( + format!("This is the vector: {}", vector_names[i]).as_str(), + )); + instructions.push(dummy_instruction); + } + + instructions[0].to_words(&mut layout.capabilities); + instructions[1].to_words(&mut layout.extensions); + instructions[2].to_words(&mut layout.ext_inst_imports); + instructions[3].to_words(&mut layout.memory_model); + instructions[4].to_words(&mut layout.entry_points); + instructions[5].to_words(&mut layout.execution_modes); + instructions[6].to_words(&mut layout.debugs); + instructions[7].to_words(&mut layout.annotations); + instructions[8].to_words(&mut layout.declarations); + instructions[9].to_words(&mut layout.function_declarations); + instructions[10].to_words(&mut layout.function_definitions); + + layout.in_words(&mut output); + + let mut index: usize = 0; + for instruction in instructions { + let wc = instruction.wc as usize; + let instruction_output = &output[index..index + wc]; + validate_instruction(instruction_output, &instruction); + index += wc; + } +} + +#[test] +fn test_instruction_set_type() { + let ty = 1; + let mut instruction = Instruction::new(Op::Constant); + assert_eq!(instruction.wc, 1); + + instruction.set_type(ty); + assert_eq!(instruction.type_id.unwrap(), ty); + assert_eq!(instruction.wc, 2); +} + +#[test] +#[should_panic] +fn test_instruction_set_type_twice() { + let ty = 1; + let mut instruction = Instruction::new(Op::Constant); + instruction.set_type(ty); + instruction.set_type(ty); +} + +#[test] +fn test_instruction_set_result() { + let result = 1; + let mut instruction = Instruction::new(Op::Constant); + assert_eq!(instruction.wc, 1); + + instruction.set_result(result); + assert_eq!(instruction.result_id.unwrap(), result); + assert_eq!(instruction.wc, 2); +} + +#[test] +#[should_panic] +fn test_instruction_set_result_twice() { + let result = 1; + let mut instruction = Instruction::new(Op::Constant); + instruction.set_result(result); + instruction.set_result(result); +} + +#[test] +fn test_instruction_add_operand() { + let operand = 1; + let mut instruction = Instruction::new(Op::Constant); + assert_eq!(instruction.operands.len(), 0); + assert_eq!(instruction.wc, 1); + + instruction.add_operand(operand); + assert_eq!(instruction.operands.len(), 1); + assert_eq!(instruction.wc, 2); +} + +#[test] +fn test_instruction_add_operands() { + let operands = vec![1, 2, 3]; + let mut instruction = Instruction::new(Op::Constant); + assert_eq!(instruction.operands.len(), 0); + assert_eq!(instruction.wc, 1); + + instruction.add_operands(operands); + assert_eq!(instruction.operands.len(), 3); + assert_eq!(instruction.wc, 4); +} + +#[test] +fn test_instruction_to_words() { + let ty = 1; + let result = 2; + let operand = 3; + let mut instruction = Instruction::new(Op::Constant); + instruction.set_type(ty); + instruction.set_result(result); + instruction.add_operand(operand); + + let mut output = vec![]; + instruction.to_words(&mut output); + validate_instruction(output.as_slice(), &instruction); +} + +fn to_word(bytes: &[u8]) -> Word { + ((bytes[0] as u32) << 16) | ((bytes[1] as u32) << 8) | bytes[2] as u32 +} diff --git a/third_party/rust/naga/src/back/spv/mod.rs b/third_party/rust/naga/src/back/spv/mod.rs new file mode 100644 index 0000000000..15f1598357 --- /dev/null +++ b/third_party/rust/naga/src/back/spv/mod.rs @@ -0,0 +1,52 @@ +mod helpers; +mod instructions; +mod layout; +mod writer; + +#[cfg(test)] +mod test_framework; + +#[cfg(test)] +mod layout_tests; + +pub use writer::Writer; + +use spirv::*; + +bitflags::bitflags! { + pub struct WriterFlags: u32 { + const NONE = 0x0; + const DEBUG = 0x1; + } +} + +struct PhysicalLayout { + magic_number: Word, + version: Word, + generator: Word, + bound: Word, + instruction_schema: Word, +} + +#[derive(Default)] +struct LogicalLayout { + capabilities: Vec<Word>, + extensions: Vec<Word>, + ext_inst_imports: Vec<Word>, + memory_model: Vec<Word>, + entry_points: Vec<Word>, + execution_modes: Vec<Word>, + debugs: Vec<Word>, + annotations: Vec<Word>, + declarations: Vec<Word>, + function_declarations: Vec<Word>, + function_definitions: Vec<Word>, +} + +pub(self) struct Instruction { + op: Op, + wc: u32, + type_id: Option<Word>, + result_id: Option<Word>, + operands: Vec<Word>, +} diff --git a/third_party/rust/naga/src/back/spv/test_framework.rs b/third_party/rust/naga/src/back/spv/test_framework.rs new file mode 100644 index 0000000000..be2fa74fe1 --- /dev/null +++ b/third_party/rust/naga/src/back/spv/test_framework.rs @@ -0,0 +1,27 @@ +pub(super) fn validate_instruction( + words: &[spirv::Word], + instruction: &crate::back::spv::Instruction, +) { + let mut inst_index = 0; + let (wc, op) = ((words[inst_index] >> 16) as u16, words[inst_index] as u16); + inst_index += 1; + + assert_eq!(wc, words.len() as u16); + assert_eq!(op, instruction.op as u16); + + if instruction.type_id.is_some() { + assert_eq!(words[inst_index], instruction.type_id.unwrap()); + inst_index += 1; + } + + if instruction.result_id.is_some() { + assert_eq!(words[inst_index], instruction.result_id.unwrap()); + inst_index += 1; + } + + let mut op_index = 0; + for i in inst_index..wc as usize { + assert_eq!(words[i as usize], instruction.operands[op_index]); + op_index += 1; + } +} diff --git a/third_party/rust/naga/src/back/spv/writer.rs b/third_party/rust/naga/src/back/spv/writer.rs new file mode 100644 index 0000000000..f1b43f5289 --- /dev/null +++ b/third_party/rust/naga/src/back/spv/writer.rs @@ -0,0 +1,1776 @@ +/*! Standard Portable Intermediate Representation (SPIR-V) backend !*/ +use super::{Instruction, LogicalLayout, PhysicalLayout, WriterFlags}; +use spirv::Word; +use std::{collections::hash_map::Entry, ops}; +use thiserror::Error; + +const BITS_PER_BYTE: crate::Bytes = 8; + +#[derive(Clone, Debug, Error)] +pub enum Error { + #[error("can't find local variable: {0:?}")] + UnknownLocalVariable(crate::LocalVariable), + #[error("bad image class for op: {0:?}")] + BadImageClass(crate::ImageClass), + #[error("not an image")] + NotImage, + #[error("empty value")] + FeatureNotImplemented(), +} + +struct Block { + label: Option<Instruction>, + body: Vec<Instruction>, + termination: Option<Instruction>, +} + +impl Block { + pub fn new() -> Self { + Block { + label: None, + body: vec![], + termination: None, + } + } +} + +struct LocalVariable { + id: Word, + name: Option<String>, + instruction: Instruction, +} + +struct Function { + signature: Option<Instruction>, + parameters: Vec<Instruction>, + variables: Vec<LocalVariable>, + blocks: Vec<Block>, +} + +impl Function { + pub fn new() -> Self { + Function { + signature: None, + parameters: vec![], + variables: vec![], + blocks: vec![], + } + } + + fn to_words(&self, sink: &mut impl Extend<Word>) { + self.signature.as_ref().unwrap().to_words(sink); + for instruction in self.parameters.iter() { + instruction.to_words(sink); + } + for (index, block) in self.blocks.iter().enumerate() { + block.label.as_ref().unwrap().to_words(sink); + if index == 0 { + for local_var in self.variables.iter() { + local_var.instruction.to_words(sink); + } + } + for instruction in block.body.iter() { + instruction.to_words(sink); + } + block.termination.as_ref().unwrap().to_words(sink); + } + } +} + +#[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)] +enum LocalType { + Void, + Scalar { + kind: crate::ScalarKind, + width: crate::Bytes, + }, + Vector { + size: crate::VectorSize, + kind: crate::ScalarKind, + width: crate::Bytes, + }, + Pointer { + base: crate::Handle<crate::Type>, + class: crate::StorageClass, + }, + SampledImage { + image_type: crate::Handle<crate::Type>, + }, +} + +#[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)] +enum LookupType { + Handle(crate::Handle<crate::Type>), + Local(LocalType), +} + +fn map_dim(dim: crate::ImageDimension) -> spirv::Dim { + match dim { + crate::ImageDimension::D1 => spirv::Dim::Dim1D, + crate::ImageDimension::D2 => spirv::Dim::Dim2D, + crate::ImageDimension::D3 => spirv::Dim::Dim2D, + crate::ImageDimension::Cube => spirv::Dim::DimCube, + } +} + +#[derive(Debug, PartialEq, Clone, Hash, Eq)] +struct LookupFunctionType { + parameter_type_ids: Vec<Word>, + return_type_id: Word, +} + +enum MaybeOwned<'a, T> { + Owned(T), + Borrowed(&'a T), +} + +impl<'a, T> ops::Deref for MaybeOwned<'a, T> { + type Target = T; + fn deref(&self) -> &T { + match *self { + MaybeOwned::Owned(ref value) => value, + MaybeOwned::Borrowed(reference) => reference, + } + } +} + +enum Dimension { + Scalar, + Vector, + Matrix, +} + +fn get_dimension(ty_inner: &crate::TypeInner) -> Dimension { + match *ty_inner { + crate::TypeInner::Scalar { .. } => Dimension::Scalar, + crate::TypeInner::Vector { .. } => Dimension::Vector, + crate::TypeInner::Matrix { .. } => Dimension::Matrix, + _ => unreachable!(), + } +} + +pub struct Writer { + physical_layout: PhysicalLayout, + logical_layout: LogicalLayout, + id_count: u32, + capabilities: crate::FastHashSet<spirv::Capability>, + debugs: Vec<Instruction>, + annotations: Vec<Instruction>, + writer_flags: WriterFlags, + void_type: Option<u32>, + lookup_type: crate::FastHashMap<LookupType, Word>, + lookup_function: crate::FastHashMap<crate::Handle<crate::Function>, Word>, + lookup_function_type: crate::FastHashMap<LookupFunctionType, Word>, + lookup_constant: crate::FastHashMap<crate::Handle<crate::Constant>, Word>, + lookup_global_variable: crate::FastHashMap<crate::Handle<crate::GlobalVariable>, Word>, +} + +// type alias, for success return of write_expression +type WriteExpressionOutput = (Word, LookupType); + +impl Writer { + pub fn new(header: &crate::Header, writer_flags: WriterFlags) -> Self { + Writer { + physical_layout: PhysicalLayout::new(header), + logical_layout: LogicalLayout::default(), + id_count: 0, + capabilities: crate::FastHashSet::default(), + debugs: vec![], + annotations: vec![], + writer_flags, + void_type: None, + lookup_type: crate::FastHashMap::default(), + lookup_function: crate::FastHashMap::default(), + lookup_function_type: crate::FastHashMap::default(), + lookup_constant: crate::FastHashMap::default(), + lookup_global_variable: crate::FastHashMap::default(), + } + } + + fn generate_id(&mut self) -> Word { + self.id_count += 1; + self.id_count + } + + fn try_add_capabilities(&mut self, capabilities: &[spirv::Capability]) { + for capability in capabilities.iter() { + self.capabilities.insert(*capability); + } + } + + fn get_type_id(&mut self, arena: &crate::Arena<crate::Type>, lookup_ty: LookupType) -> Word { + if let Entry::Occupied(e) = self.lookup_type.entry(lookup_ty) { + *e.get() + } else { + match lookup_ty { + LookupType::Handle(handle) => match arena[handle].inner { + crate::TypeInner::Scalar { kind, width } => self + .get_type_id(arena, LookupType::Local(LocalType::Scalar { kind, width })), + _ => self.write_type_declaration_arena(arena, handle), + }, + LookupType::Local(local_ty) => self.write_type_declaration_local(arena, local_ty), + } + } + } + + fn get_constant_id( + &mut self, + handle: crate::Handle<crate::Constant>, + ir_module: &crate::Module, + ) -> Word { + match self.lookup_constant.entry(handle) { + Entry::Occupied(e) => *e.get(), + _ => { + let (instruction, id) = self.write_constant_type(handle, ir_module); + instruction.to_words(&mut self.logical_layout.declarations); + id + } + } + } + + fn get_global_variable_id( + &mut self, + ir_module: &crate::Module, + handle: crate::Handle<crate::GlobalVariable>, + ) -> Word { + match self.lookup_global_variable.entry(handle) { + Entry::Occupied(e) => *e.get(), + _ => { + let (instruction, id) = self.write_global_variable(ir_module, handle); + instruction.to_words(&mut self.logical_layout.declarations); + id + } + } + } + + fn get_function_return_type( + &mut self, + ty: Option<crate::Handle<crate::Type>>, + arena: &crate::Arena<crate::Type>, + ) -> Word { + match ty { + Some(handle) => self.get_type_id(arena, LookupType::Handle(handle)), + None => match self.void_type { + Some(id) => id, + None => { + let id = self.generate_id(); + self.void_type = Some(id); + super::instructions::instruction_type_void(id) + .to_words(&mut self.logical_layout.declarations); + id + } + }, + } + } + + fn get_pointer_id( + &mut self, + arena: &crate::Arena<crate::Type>, + handle: crate::Handle<crate::Type>, + class: crate::StorageClass, + ) -> Word { + let ty = &arena[handle]; + let ty_id = self.get_type_id(arena, LookupType::Handle(handle)); + match ty.inner { + crate::TypeInner::Pointer { .. } => ty_id, + _ => { + match self + .lookup_type + .entry(LookupType::Local(LocalType::Pointer { + base: handle, + class, + })) { + Entry::Occupied(e) => *e.get(), + _ => { + let id = + self.create_pointer(ty_id, self.parse_to_spirv_storage_class(class)); + self.lookup_type.insert( + LookupType::Local(LocalType::Pointer { + base: handle, + class, + }), + id, + ); + id + } + } + } + } + } + + fn create_pointer(&mut self, ty_id: Word, class: spirv::StorageClass) -> Word { + let id = self.generate_id(); + let instruction = super::instructions::instruction_type_pointer(id, class, ty_id); + instruction.to_words(&mut self.logical_layout.declarations); + id + } + + fn create_constant(&mut self, type_id: Word, value: &[Word]) -> Word { + let id = self.generate_id(); + let instruction = super::instructions::instruction_constant(type_id, id, value); + instruction.to_words(&mut self.logical_layout.declarations); + id + } + + fn write_function( + &mut self, + ir_function: &crate::Function, + ir_module: &crate::Module, + ) -> spirv::Word { + let mut function = Function::new(); + + for (_, variable) in ir_function.local_variables.iter() { + let id = self.generate_id(); + + let init_word = variable + .init + .map(|constant| self.get_constant_id(constant, ir_module)); + + let pointer_id = + self.get_pointer_id(&ir_module.types, variable.ty, crate::StorageClass::Function); + function.variables.push(LocalVariable { + id, + name: variable.name.clone(), + instruction: super::instructions::instruction_variable( + pointer_id, + id, + spirv::StorageClass::Function, + init_word, + ), + }); + } + + let return_type_id = + self.get_function_return_type(ir_function.return_type, &ir_module.types); + let mut parameter_type_ids = Vec::with_capacity(ir_function.arguments.len()); + + let mut function_parameter_pointer_ids = vec![]; + + for argument in ir_function.arguments.iter() { + let id = self.generate_id(); + let pointer_id = + self.get_pointer_id(&ir_module.types, argument.ty, crate::StorageClass::Function); + + function_parameter_pointer_ids.push(pointer_id); + parameter_type_ids + .push(self.get_type_id(&ir_module.types, LookupType::Handle(argument.ty))); + function + .parameters + .push(super::instructions::instruction_function_parameter( + pointer_id, id, + )); + } + + let lookup_function_type = LookupFunctionType { + return_type_id, + parameter_type_ids, + }; + + let function_id = self.generate_id(); + let function_type = + self.get_function_type(lookup_function_type, function_parameter_pointer_ids); + function.signature = Some(super::instructions::instruction_function( + return_type_id, + function_id, + spirv::FunctionControl::empty(), + function_type, + )); + + self.write_block(&ir_function.body, ir_module, ir_function, &mut function); + + function.to_words(&mut self.logical_layout.function_definitions); + super::instructions::instruction_function_end() + .to_words(&mut self.logical_layout.function_definitions); + + function_id + } + + // TODO Move to instructions module + fn write_entry_point( + &mut self, + entry_point: &crate::EntryPoint, + stage: crate::ShaderStage, + name: &str, + ir_module: &crate::Module, + ) -> Instruction { + let function_id = self.write_function(&entry_point.function, ir_module); + + let exec_model = match stage { + crate::ShaderStage::Vertex => spirv::ExecutionModel::Vertex, + crate::ShaderStage::Fragment { .. } => spirv::ExecutionModel::Fragment, + crate::ShaderStage::Compute { .. } => spirv::ExecutionModel::GLCompute, + }; + + let mut interface_ids = vec![]; + for ((handle, _), &usage) in ir_module + .global_variables + .iter() + .filter(|&(_, var)| { + var.class == crate::StorageClass::Input || var.class == crate::StorageClass::Output + }) + .zip(&entry_point.function.global_usage) + { + if usage.contains(crate::GlobalUse::STORE) || usage.contains(crate::GlobalUse::LOAD) { + let id = self.get_global_variable_id(ir_module, handle); + interface_ids.push(id); + } + } + + self.try_add_capabilities(exec_model.required_capabilities()); + match stage { + crate::ShaderStage::Vertex => {} + crate::ShaderStage::Fragment => { + let execution_mode = spirv::ExecutionMode::OriginUpperLeft; + self.try_add_capabilities(execution_mode.required_capabilities()); + super::instructions::instruction_execution_mode(function_id, execution_mode) + .to_words(&mut self.logical_layout.execution_modes); + } + crate::ShaderStage::Compute => {} + } + + if self.writer_flags.contains(WriterFlags::DEBUG) { + self.debugs + .push(super::instructions::instruction_name(function_id, name)); + } + + super::instructions::instruction_entry_point( + exec_model, + function_id, + name, + interface_ids.as_slice(), + ) + } + + fn write_scalar(&self, id: Word, kind: crate::ScalarKind, width: crate::Bytes) -> Instruction { + let bits = (width * BITS_PER_BYTE) as u32; + match kind { + crate::ScalarKind::Sint => super::instructions::instruction_type_int( + id, + bits, + super::instructions::Signedness::Signed, + ), + crate::ScalarKind::Uint => super::instructions::instruction_type_int( + id, + bits, + super::instructions::Signedness::Unsigned, + ), + crate::ScalarKind::Float => super::instructions::instruction_type_float(id, bits), + crate::ScalarKind::Bool => super::instructions::instruction_type_bool(id), + } + } + + fn parse_to_spirv_storage_class(&self, class: crate::StorageClass) -> spirv::StorageClass { + match class { + crate::StorageClass::Handle => spirv::StorageClass::UniformConstant, + crate::StorageClass::Function => spirv::StorageClass::Function, + crate::StorageClass::Input => spirv::StorageClass::Input, + crate::StorageClass::Output => spirv::StorageClass::Output, + crate::StorageClass::Private => spirv::StorageClass::Private, + crate::StorageClass::Storage if self.physical_layout.supports_storage_buffers() => { + spirv::StorageClass::StorageBuffer + } + crate::StorageClass::Storage | crate::StorageClass::Uniform => { + spirv::StorageClass::Uniform + } + crate::StorageClass::WorkGroup => spirv::StorageClass::Workgroup, + crate::StorageClass::PushConstant => spirv::StorageClass::PushConstant, + } + } + + fn write_type_declaration_local( + &mut self, + arena: &crate::Arena<crate::Type>, + local_ty: LocalType, + ) -> Word { + let id = self.generate_id(); + let instruction = match local_ty { + LocalType::Void => unreachable!(), + LocalType::Scalar { kind, width } => self.write_scalar(id, kind, width), + LocalType::Vector { size, kind, width } => { + let scalar_id = + self.get_type_id(arena, LookupType::Local(LocalType::Scalar { kind, width })); + super::instructions::instruction_type_vector(id, scalar_id, size) + } + LocalType::Pointer { .. } => unimplemented!(), + LocalType::SampledImage { image_type } => { + let image_type_id = self.get_type_id(arena, LookupType::Handle(image_type)); + super::instructions::instruction_type_sampled_image(id, image_type_id) + } + }; + + self.lookup_type.insert(LookupType::Local(local_ty), id); + instruction.to_words(&mut self.logical_layout.declarations); + id + } + + fn write_type_declaration_arena( + &mut self, + arena: &crate::Arena<crate::Type>, + handle: crate::Handle<crate::Type>, + ) -> Word { + let ty = &arena[handle]; + let id = self.generate_id(); + + let instruction = match ty.inner { + crate::TypeInner::Scalar { kind, width } => { + self.lookup_type + .insert(LookupType::Local(LocalType::Scalar { kind, width }), id); + self.write_scalar(id, kind, width) + } + crate::TypeInner::Vector { size, kind, width } => { + let scalar_id = + self.get_type_id(arena, LookupType::Local(LocalType::Scalar { kind, width })); + self.lookup_type.insert( + LookupType::Local(LocalType::Vector { size, kind, width }), + id, + ); + super::instructions::instruction_type_vector(id, scalar_id, size) + } + crate::TypeInner::Matrix { + columns, + rows: _, + width, + } => { + let vector_id = self.get_type_id( + arena, + LookupType::Local(LocalType::Vector { + size: columns, + kind: crate::ScalarKind::Float, + width, + }), + ); + super::instructions::instruction_type_matrix(id, vector_id, columns) + } + crate::TypeInner::Image { + dim, + arrayed, + class, + } => { + let width = 4; + let local_type = match class { + crate::ImageClass::Sampled { kind, multi: _ } => { + LocalType::Scalar { kind, width } + } + crate::ImageClass::Depth => LocalType::Scalar { + kind: crate::ScalarKind::Float, + width, + }, + crate::ImageClass::Storage(format) => LocalType::Scalar { + kind: format.into(), + width, + }, + }; + let type_id = self.get_type_id(arena, LookupType::Local(local_type)); + let dim = map_dim(dim); + self.try_add_capabilities(dim.required_capabilities()); + super::instructions::instruction_type_image(id, type_id, dim, arrayed, class) + } + crate::TypeInner::Sampler { comparison: _ } => { + super::instructions::instruction_type_sampler(id) + } + crate::TypeInner::Array { base, size, stride } => { + if let Some(array_stride) = stride { + self.annotations + .push(super::instructions::instruction_decorate( + id, + spirv::Decoration::ArrayStride, + &[array_stride.get()], + )); + } + + let type_id = self.get_type_id(arena, LookupType::Handle(base)); + match size { + crate::ArraySize::Constant(const_handle) => { + let length_id = self.lookup_constant[&const_handle]; + super::instructions::instruction_type_array(id, type_id, length_id) + } + crate::ArraySize::Dynamic => { + super::instructions::instruction_type_runtime_array(id, type_id) + } + } + } + crate::TypeInner::Struct { ref members } => { + let mut member_ids = Vec::with_capacity(members.len()); + for member in members { + let member_id = self.get_type_id(arena, LookupType::Handle(member.ty)); + member_ids.push(member_id); + } + super::instructions::instruction_type_struct(id, member_ids.as_slice()) + } + crate::TypeInner::Pointer { base, class } => { + let type_id = self.get_type_id(arena, LookupType::Handle(base)); + self.lookup_type + .insert(LookupType::Local(LocalType::Pointer { base, class }), id); + super::instructions::instruction_type_pointer( + id, + self.parse_to_spirv_storage_class(class), + type_id, + ) + } + }; + + self.lookup_type.insert(LookupType::Handle(handle), id); + instruction.to_words(&mut self.logical_layout.declarations); + id + } + + fn write_constant_type( + &mut self, + handle: crate::Handle<crate::Constant>, + ir_module: &crate::Module, + ) -> (Instruction, Word) { + let id = self.generate_id(); + self.lookup_constant.insert(handle, id); + let constant = &ir_module.constants[handle]; + let arena = &ir_module.types; + + match constant.inner { + crate::ConstantInner::Sint(val) => { + let ty = &ir_module.types[constant.ty]; + let type_id = self.get_type_id(arena, LookupType::Handle(constant.ty)); + + let instruction = match ty.inner { + crate::TypeInner::Scalar { kind: _, width } => match width { + 4 => super::instructions::instruction_constant(type_id, id, &[val as u32]), + 8 => { + let (low, high) = ((val >> 32) as u32, val as u32); + super::instructions::instruction_constant(type_id, id, &[low, high]) + } + _ => unreachable!(), + }, + _ => unreachable!(), + }; + (instruction, id) + } + crate::ConstantInner::Uint(val) => { + let ty = &ir_module.types[constant.ty]; + let type_id = self.get_type_id(arena, LookupType::Handle(constant.ty)); + + let instruction = match ty.inner { + crate::TypeInner::Scalar { kind: _, width } => match width { + 4 => super::instructions::instruction_constant(type_id, id, &[val as u32]), + 8 => { + let (low, high) = ((val >> 32) as u32, val as u32); + super::instructions::instruction_constant(type_id, id, &[low, high]) + } + _ => unreachable!(), + }, + _ => unreachable!(), + }; + + (instruction, id) + } + crate::ConstantInner::Float(val) => { + let ty = &ir_module.types[constant.ty]; + let type_id = self.get_type_id(arena, LookupType::Handle(constant.ty)); + + let instruction = match ty.inner { + crate::TypeInner::Scalar { kind: _, width } => match width { + 4 => super::instructions::instruction_constant( + type_id, + id, + &[(val as f32).to_bits()], + ), + 8 => { + let bits = f64::to_bits(val); + let (low, high) = ((bits >> 32) as u32, bits as u32); + super::instructions::instruction_constant(type_id, id, &[low, high]) + } + _ => unreachable!(), + }, + _ => unreachable!(), + }; + (instruction, id) + } + crate::ConstantInner::Bool(val) => { + let type_id = self.get_type_id(arena, LookupType::Handle(constant.ty)); + + let instruction = if val { + super::instructions::instruction_constant_true(type_id, id) + } else { + super::instructions::instruction_constant_false(type_id, id) + }; + + (instruction, id) + } + crate::ConstantInner::Composite(ref constituents) => { + let mut constituent_ids = Vec::with_capacity(constituents.len()); + for constituent in constituents.iter() { + let constituent_id = self.get_constant_id(*constituent, &ir_module); + constituent_ids.push(constituent_id); + } + + let type_id = self.get_type_id(arena, LookupType::Handle(constant.ty)); + let instruction = super::instructions::instruction_constant_composite( + type_id, + id, + constituent_ids.as_slice(), + ); + (instruction, id) + } + } + } + + fn write_global_variable( + &mut self, + ir_module: &crate::Module, + handle: crate::Handle<crate::GlobalVariable>, + ) -> (Instruction, Word) { + let global_variable = &ir_module.global_variables[handle]; + let id = self.generate_id(); + + let class = self.parse_to_spirv_storage_class(global_variable.class); + self.try_add_capabilities(class.required_capabilities()); + + let init_word = global_variable + .init + .map(|constant| self.get_constant_id(constant, ir_module)); + let pointer_id = + self.get_pointer_id(&ir_module.types, global_variable.ty, global_variable.class); + let instruction = + super::instructions::instruction_variable(pointer_id, id, class, init_word); + + if self.writer_flags.contains(WriterFlags::DEBUG) { + if let Some(ref name) = global_variable.name { + self.debugs + .push(super::instructions::instruction_name(id, name.as_str())); + } + } + + if let Some(interpolation) = global_variable.interpolation { + let decoration = match interpolation { + crate::Interpolation::Linear => Some(spirv::Decoration::NoPerspective), + crate::Interpolation::Flat => Some(spirv::Decoration::Flat), + crate::Interpolation::Patch => Some(spirv::Decoration::Patch), + crate::Interpolation::Centroid => Some(spirv::Decoration::Centroid), + crate::Interpolation::Sample => Some(spirv::Decoration::Sample), + crate::Interpolation::Perspective => None, + }; + if let Some(decoration) = decoration { + self.annotations + .push(super::instructions::instruction_decorate( + id, + decoration, + &[], + )); + } + } + + match *global_variable.binding.as_ref().unwrap() { + crate::Binding::Location(location) => { + self.annotations + .push(super::instructions::instruction_decorate( + id, + spirv::Decoration::Location, + &[location], + )); + } + crate::Binding::Resource { group, binding } => { + self.annotations + .push(super::instructions::instruction_decorate( + id, + spirv::Decoration::DescriptorSet, + &[group], + )); + self.annotations + .push(super::instructions::instruction_decorate( + id, + spirv::Decoration::Binding, + &[binding], + )); + } + crate::Binding::BuiltIn(built_in) => { + let built_in = match built_in { + crate::BuiltIn::BaseInstance => spirv::BuiltIn::BaseInstance, + crate::BuiltIn::BaseVertex => spirv::BuiltIn::BaseVertex, + crate::BuiltIn::ClipDistance => spirv::BuiltIn::ClipDistance, + crate::BuiltIn::InstanceIndex => spirv::BuiltIn::InstanceIndex, + crate::BuiltIn::Position => spirv::BuiltIn::Position, + crate::BuiltIn::VertexIndex => spirv::BuiltIn::VertexIndex, + crate::BuiltIn::PointSize => spirv::BuiltIn::PointSize, + crate::BuiltIn::FragCoord => spirv::BuiltIn::FragCoord, + crate::BuiltIn::FrontFacing => spirv::BuiltIn::FrontFacing, + crate::BuiltIn::SampleIndex => spirv::BuiltIn::SampleId, + crate::BuiltIn::FragDepth => spirv::BuiltIn::FragDepth, + crate::BuiltIn::GlobalInvocationId => spirv::BuiltIn::GlobalInvocationId, + crate::BuiltIn::LocalInvocationId => spirv::BuiltIn::LocalInvocationId, + crate::BuiltIn::LocalInvocationIndex => spirv::BuiltIn::LocalInvocationIndex, + crate::BuiltIn::WorkGroupId => spirv::BuiltIn::WorkgroupId, + }; + + self.annotations + .push(super::instructions::instruction_decorate( + id, + spirv::Decoration::BuiltIn, + &[built_in as u32], + )); + } + } + + // TODO Initializer is optional and not (yet) included in the IR + + self.lookup_global_variable.insert(handle, id); + (instruction, id) + } + + fn get_function_type( + &mut self, + lookup_function_type: LookupFunctionType, + parameter_pointer_ids: Vec<Word>, + ) -> Word { + match self + .lookup_function_type + .entry(lookup_function_type.clone()) + { + Entry::Occupied(e) => *e.get(), + _ => { + let id = self.generate_id(); + let instruction = super::instructions::instruction_type_function( + id, + lookup_function_type.return_type_id, + parameter_pointer_ids.as_slice(), + ); + instruction.to_words(&mut self.logical_layout.declarations); + self.lookup_function_type.insert(lookup_function_type, id); + id + } + } + } + + fn write_composite_construct( + &mut self, + base_type_id: Word, + constituent_ids: &[Word], + block: &mut Block, + ) -> Word { + let id = self.generate_id(); + block + .body + .push(super::instructions::instruction_composite_construct( + base_type_id, + id, + constituent_ids, + )); + id + } + + fn get_type_inner<'a>( + &self, + ty_arena: &'a crate::Arena<crate::Type>, + lookup_ty: LookupType, + ) -> MaybeOwned<'a, crate::TypeInner> { + match lookup_ty { + LookupType::Handle(handle) => MaybeOwned::Borrowed(&ty_arena[handle].inner), + LookupType::Local(local_ty) => match local_ty { + LocalType::Scalar { kind, width } => { + MaybeOwned::Owned(crate::TypeInner::Scalar { kind, width }) + } + LocalType::Vector { size, kind, width } => { + MaybeOwned::Owned(crate::TypeInner::Vector { size, kind, width }) + } + LocalType::Pointer { base, class } => { + MaybeOwned::Owned(crate::TypeInner::Pointer { base, class }) + } + _ => unreachable!(), + }, + } + } + + fn write_expression<'a>( + &mut self, + ir_module: &'a crate::Module, + ir_function: &crate::Function, + expression: &crate::Expression, + block: &mut Block, + function: &mut Function, + ) -> Result<WriteExpressionOutput, Error> { + match *expression { + crate::Expression::Access { base, index } => { + let id = self.generate_id(); + + let (base_id, base_lookup_ty) = self.write_expression( + ir_module, + ir_function, + &ir_function.expressions[base], + block, + function, + )?; + let (index_id, _) = self.write_expression( + ir_module, + ir_function, + &ir_function.expressions[index], + block, + function, + )?; + + let base_ty_inner = self.get_type_inner(&ir_module.types, base_lookup_ty); + + let (pointer_id, type_id, lookup_ty) = match *base_ty_inner { + crate::TypeInner::Vector { kind, width, .. } => { + let scalar_id = self.get_type_id( + &ir_module.types, + LookupType::Local(LocalType::Scalar { kind, width }), + ); + ( + self.create_pointer(scalar_id, spirv::StorageClass::Function), + scalar_id, + LookupType::Local(LocalType::Scalar { kind, width }), + ) + } + _ => unimplemented!(), + }; + + block + .body + .push(super::instructions::instruction_access_chain( + pointer_id, + id, + base_id, + &[index_id], + )); + + let load_id = self.generate_id(); + block.body.push(super::instructions::instruction_load( + type_id, load_id, id, None, + )); + + Ok((load_id, lookup_ty)) + } + crate::Expression::AccessIndex { base, index } => { + let id = self.generate_id(); + let (base_id, base_lookup_ty) = self + .write_expression( + ir_module, + ir_function, + &ir_function.expressions[base], + block, + function, + ) + .unwrap(); + + let base_ty_inner = self.get_type_inner(&ir_module.types, base_lookup_ty); + + let (pointer_id, type_id, lookup_ty) = match *base_ty_inner { + crate::TypeInner::Vector { kind, width, .. } => { + let scalar_id = self.get_type_id( + &ir_module.types, + LookupType::Local(LocalType::Scalar { kind, width }), + ); + ( + self.create_pointer(scalar_id, spirv::StorageClass::Function), + scalar_id, + LookupType::Local(LocalType::Scalar { kind, width }), + ) + } + crate::TypeInner::Struct { ref members } => { + let member = &members[index as usize]; + let type_id = + self.get_type_id(&ir_module.types, LookupType::Handle(member.ty)); + ( + self.create_pointer(type_id, spirv::StorageClass::Uniform), + type_id, + LookupType::Handle(member.ty), + ) + } + _ => unimplemented!(), + }; + + let const_ty_id = self.get_type_id( + &ir_module.types, + LookupType::Local(LocalType::Scalar { + kind: crate::ScalarKind::Sint, + width: 4, + }), + ); + let const_id = self.create_constant(const_ty_id, &[index]); + + block + .body + .push(super::instructions::instruction_access_chain( + pointer_id, + id, + base_id, + &[const_id], + )); + + let load_id = self.generate_id(); + block.body.push(super::instructions::instruction_load( + type_id, load_id, id, None, + )); + + Ok((load_id, lookup_ty)) + } + crate::Expression::GlobalVariable(handle) => { + let var = &ir_module.global_variables[handle]; + let id = self.get_global_variable_id(&ir_module, handle); + + Ok((id, LookupType::Handle(var.ty))) + } + crate::Expression::Constant(handle) => { + let var = &ir_module.constants[handle]; + let id = self.get_constant_id(handle, ir_module); + Ok((id, LookupType::Handle(var.ty))) + } + crate::Expression::Compose { ty, ref components } => { + let base_type_id = self.get_type_id(&ir_module.types, LookupType::Handle(ty)); + + let mut constituent_ids = Vec::with_capacity(components.len()); + for component in components { + let expression = &ir_function.expressions[*component]; + let (component_id, component_local_ty) = self.write_expression( + ir_module, + &ir_function, + expression, + block, + function, + )?; + + let component_id = match expression { + crate::Expression::LocalVariable(_) + | crate::Expression::GlobalVariable(_) => { + let load_id = self.generate_id(); + block.body.push(super::instructions::instruction_load( + self.get_type_id(&ir_module.types, component_local_ty), + load_id, + component_id, + None, + )); + load_id + } + _ => component_id, + }; + + constituent_ids.push(component_id); + } + let constituent_ids_slice = constituent_ids.as_slice(); + + let id = match ir_module.types[ty].inner { + crate::TypeInner::Vector { .. } => { + self.write_composite_construct(base_type_id, constituent_ids_slice, block) + } + crate::TypeInner::Matrix { + rows, + columns, + width, + } => { + let vector_type_id = self.get_type_id( + &ir_module.types, + LookupType::Local(LocalType::Vector { + width, + kind: crate::ScalarKind::Float, + size: columns, + }), + ); + + let capacity = match rows { + crate::VectorSize::Bi => 2, + crate::VectorSize::Tri => 3, + crate::VectorSize::Quad => 4, + }; + + let mut vector_ids = Vec::with_capacity(capacity); + + for _ in 0..capacity { + let vector_id = self.write_composite_construct( + vector_type_id, + constituent_ids_slice, + block, + ); + vector_ids.push(vector_id); + } + + self.write_composite_construct(base_type_id, vector_ids.as_slice(), block) + } + _ => unreachable!(), + }; + + Ok((id, LookupType::Handle(ty))) + } + crate::Expression::Binary { op, left, right } => { + let id = self.generate_id(); + let left_expression = &ir_function.expressions[left]; + let right_expression = &ir_function.expressions[right]; + let (left_id, left_lookup_ty) = self.write_expression( + ir_module, + ir_function, + left_expression, + block, + function, + )?; + let (right_id, right_lookup_ty) = self.write_expression( + ir_module, + ir_function, + right_expression, + block, + function, + )?; + + let left_lookup_ty = left_lookup_ty; + let right_lookup_ty = right_lookup_ty; + + let left_ty_inner = self.get_type_inner(&ir_module.types, left_lookup_ty); + let right_ty_inner = self.get_type_inner(&ir_module.types, right_lookup_ty); + + let left_result_type_id = self.get_type_id(&ir_module.types, left_lookup_ty); + + let right_result_type_id = self.get_type_id(&ir_module.types, right_lookup_ty); + + let left_id = match *left_expression { + crate::Expression::LocalVariable(_) | crate::Expression::GlobalVariable(_) => { + let load_id = self.generate_id(); + block.body.push(super::instructions::instruction_load( + left_result_type_id, + load_id, + left_id, + None, + )); + load_id + } + _ => left_id, + }; + + let right_id = match *right_expression { + crate::Expression::LocalVariable(..) + | crate::Expression::GlobalVariable(..) => { + let load_id = self.generate_id(); + block.body.push(super::instructions::instruction_load( + right_result_type_id, + load_id, + right_id, + None, + )); + load_id + } + _ => right_id, + }; + + let left_dimension = get_dimension(&left_ty_inner); + let right_dimension = get_dimension(&right_ty_inner); + + let (instruction, lookup_ty) = match op { + crate::BinaryOperator::Multiply => match (left_dimension, right_dimension) { + (Dimension::Vector, Dimension::Scalar { .. }) => ( + super::instructions::instruction_vector_times_scalar( + left_result_type_id, + id, + left_id, + right_id, + ), + left_lookup_ty, + ), + (Dimension::Vector, Dimension::Matrix) => ( + super::instructions::instruction_vector_times_matrix( + left_result_type_id, + id, + left_id, + right_id, + ), + left_lookup_ty, + ), + (Dimension::Matrix, Dimension::Scalar { .. }) => ( + super::instructions::instruction_matrix_times_scalar( + left_result_type_id, + id, + left_id, + right_id, + ), + left_lookup_ty, + ), + (Dimension::Matrix, Dimension::Vector) => ( + super::instructions::instruction_matrix_times_vector( + right_result_type_id, + id, + left_id, + right_id, + ), + right_lookup_ty, + ), + (Dimension::Matrix, Dimension::Matrix) => ( + super::instructions::instruction_matrix_times_matrix( + left_result_type_id, + id, + left_id, + right_id, + ), + left_lookup_ty, + ), + (Dimension::Vector, Dimension::Vector) + | (Dimension::Scalar, Dimension::Scalar) + if left_ty_inner.scalar_kind() == Some(crate::ScalarKind::Float) => + { + ( + super::instructions::instruction_f_mul( + left_result_type_id, + id, + left_id, + right_id, + ), + left_lookup_ty, + ) + } + (Dimension::Vector, Dimension::Vector) + | (Dimension::Scalar, Dimension::Scalar) => ( + super::instructions::instruction_i_mul( + left_result_type_id, + id, + left_id, + right_id, + ), + left_lookup_ty, + ), + _ => unreachable!(), + }, + crate::BinaryOperator::Subtract => match *left_ty_inner { + crate::TypeInner::Scalar { kind, .. } => match kind { + crate::ScalarKind::Sint | crate::ScalarKind::Uint => ( + super::instructions::instruction_i_sub( + left_result_type_id, + id, + left_id, + right_id, + ), + left_lookup_ty, + ), + crate::ScalarKind::Float => ( + super::instructions::instruction_f_sub( + left_result_type_id, + id, + left_id, + right_id, + ), + left_lookup_ty, + ), + _ => unreachable!(), + }, + _ => unreachable!(), + }, + crate::BinaryOperator::And => ( + super::instructions::instruction_bitwise_and( + left_result_type_id, + id, + left_id, + right_id, + ), + left_lookup_ty, + ), + _ => unimplemented!("{:?}", op), + }; + + block.body.push(instruction); + Ok((id, lookup_ty)) + } + crate::Expression::LocalVariable(variable) => { + let var = &ir_function.local_variables[variable]; + function + .variables + .iter() + .find(|&v| v.name.as_ref().unwrap() == var.name.as_ref().unwrap()) + .map(|local_var| (local_var.id, LookupType::Handle(var.ty))) + .ok_or_else(|| Error::UnknownLocalVariable(var.clone())) + } + crate::Expression::FunctionArgument(index) => { + let handle = ir_function.arguments[index as usize].ty; + let type_id = self.get_type_id(&ir_module.types, LookupType::Handle(handle)); + let load_id = self.generate_id(); + + block.body.push(super::instructions::instruction_load( + type_id, + load_id, + function.parameters[index as usize].result_id.unwrap(), + None, + )); + Ok((load_id, LookupType::Handle(handle))) + } + crate::Expression::Call { + ref origin, + ref arguments, + } => match *origin { + crate::FunctionOrigin::Local(local_function) => { + let origin_function = &ir_module.functions[local_function]; + let id = self.generate_id(); + let mut argument_ids = vec![]; + + for argument in arguments { + let expression = &ir_function.expressions[*argument]; + let (id, lookup_ty) = self.write_expression( + ir_module, + ir_function, + expression, + block, + function, + )?; + + // Create variable - OpVariable + // Store value to variable - OpStore + // Use id of variable + + let handle = match lookup_ty { + LookupType::Handle(handle) => handle, + LookupType::Local(_) => unreachable!(), + }; + + let pointer_id = self.get_pointer_id( + &ir_module.types, + handle, + crate::StorageClass::Function, + ); + + let variable_id = self.generate_id(); + function.variables.push(LocalVariable { + id: variable_id, + name: None, + instruction: super::instructions::instruction_variable( + pointer_id, + variable_id, + spirv::StorageClass::Function, + None, + ), + }); + block.body.push(super::instructions::instruction_store( + variable_id, + id, + None, + )); + argument_ids.push(variable_id); + } + + let return_type_id = self + .get_function_return_type(origin_function.return_type, &ir_module.types); + + block + .body + .push(super::instructions::instruction_function_call( + return_type_id, + id, + *self.lookup_function.get(&local_function).unwrap(), + argument_ids.as_slice(), + )); + + let result_type = match origin_function.return_type { + Some(ty_handle) => LookupType::Handle(ty_handle), + None => LookupType::Local(LocalType::Void), + }; + Ok((id, result_type)) + } + _ => unimplemented!("{:?}", origin), + }, + crate::Expression::As { + expr, + kind, + convert, + } => { + if !convert { + return Err(Error::FeatureNotImplemented()); + } + + let (expr_id, expr_type) = self.write_expression( + ir_module, + ir_function, + &ir_function.expressions[expr], + block, + function, + )?; + + let expr_type_inner = self.get_type_inner(&ir_module.types, expr_type); + + let (expr_kind, local_type) = match *expr_type_inner { + crate::TypeInner::Scalar { + kind: expr_kind, + width, + } => (expr_kind, LocalType::Scalar { kind, width }), + crate::TypeInner::Vector { + size, + kind: expr_kind, + width, + } => (expr_kind, LocalType::Vector { size, kind, width }), + _ => unreachable!(), + }; + + let lookup_type = LookupType::Local(local_type); + let op = match (expr_kind, kind) { + _ if !convert => spirv::Op::Bitcast, + (crate::ScalarKind::Float, crate::ScalarKind::Uint) => spirv::Op::ConvertFToU, + (crate::ScalarKind::Float, crate::ScalarKind::Sint) => spirv::Op::ConvertFToS, + (crate::ScalarKind::Sint, crate::ScalarKind::Float) => spirv::Op::ConvertSToF, + (crate::ScalarKind::Uint, crate::ScalarKind::Float) => spirv::Op::ConvertUToF, + // We assume it's either an identity cast, or int-uint. + // In both cases no SPIR-V instructions need to be generated. + _ => { + let id = match ir_function.expressions[expr] { + crate::Expression::LocalVariable(_) + | crate::Expression::GlobalVariable(_) => { + let load_id = self.generate_id(); + let kind_type_id = self.get_type_id(&ir_module.types, expr_type); + block.body.push(super::instructions::instruction_load( + kind_type_id, + load_id, + expr_id, + None, + )); + load_id + } + _ => expr_id, + }; + return Ok((id, lookup_type)); + } + }; + + let id = self.generate_id(); + let kind_type_id = self.get_type_id(&ir_module.types, lookup_type); + let instruction = + super::instructions::instruction_unary(op, kind_type_id, id, expr_id); + block.body.push(instruction); + + Ok((id, lookup_type)) + } + crate::Expression::ImageSample { + image, + sampler, + coordinate, + level: _, + depth_ref: _, + } => { + // image + let image_expression = &ir_function.expressions[image]; + let (image_id, image_lookup_ty) = self.write_expression( + ir_module, + ir_function, + image_expression, + block, + function, + )?; + + let image_result_type_id = self.get_type_id(&ir_module.types, image_lookup_ty); + let image_id = match *image_expression { + crate::Expression::LocalVariable(_) | crate::Expression::GlobalVariable(_) => { + let load_id = self.generate_id(); + block.body.push(super::instructions::instruction_load( + image_result_type_id, + load_id, + image_id, + None, + )); + load_id + } + _ => image_id, + }; + + let image_ty = match image_lookup_ty { + LookupType::Handle(handle) => handle, + LookupType::Local(_) => unreachable!(), + }; + + // OpTypeSampledImage + let sampled_image_type_id = self.get_type_id( + &ir_module.types, + LookupType::Local(LocalType::SampledImage { + image_type: image_ty, + }), + ); + + // sampler + let sampler_expression = &ir_function.expressions[sampler]; + let (sampler_id, sampler_lookup_ty) = self.write_expression( + ir_module, + ir_function, + sampler_expression, + block, + function, + )?; + + let sampler_result_type_id = self.get_type_id(&ir_module.types, sampler_lookup_ty); + let sampler_id = match *sampler_expression { + crate::Expression::LocalVariable(_) | crate::Expression::GlobalVariable(_) => { + let load_id = self.generate_id(); + block.body.push(super::instructions::instruction_load( + sampler_result_type_id, + load_id, + sampler_id, + None, + )); + load_id + } + _ => sampler_id, + }; + + // coordinate + let coordinate_expression = &ir_function.expressions[coordinate]; + let (coordinate_id, coordinate_lookup_ty) = self.write_expression( + ir_module, + ir_function, + coordinate_expression, + block, + function, + )?; + + let coordinate_result_type_id = + self.get_type_id(&ir_module.types, coordinate_lookup_ty); + let coordinate_id = match *coordinate_expression { + crate::Expression::LocalVariable(_) | crate::Expression::GlobalVariable(_) => { + let load_id = self.generate_id(); + block.body.push(super::instructions::instruction_load( + coordinate_result_type_id, + load_id, + coordinate_id, + None, + )); + load_id + } + _ => coordinate_id, + }; + + // component kind + let image_type = &ir_module.types[image_ty]; + let image_sample_result_type = + if let crate::TypeInner::Image { class, .. } = image_type.inner { + let width = 4; + LookupType::Local(match class { + crate::ImageClass::Sampled { kind, multi: _ } => LocalType::Vector { + kind, + width, + size: crate::VectorSize::Quad, + }, + crate::ImageClass::Depth => LocalType::Scalar { + kind: crate::ScalarKind::Float, + width, + }, + _ => return Err(Error::BadImageClass(class)), + }) + } else { + return Err(Error::NotImage); + }; + + let sampled_image_id = self.generate_id(); + block + .body + .push(super::instructions::instruction_sampled_image( + sampled_image_type_id, + sampled_image_id, + image_id, + sampler_id, + )); + let id = self.generate_id(); + let image_sample_result_type_id = + self.get_type_id(&ir_module.types, image_sample_result_type); + block + .body + .push(super::instructions::instruction_image_sample_implicit_lod( + image_sample_result_type_id, + id, + sampled_image_id, + coordinate_id, + )); + Ok((id, image_sample_result_type)) + } + _ => unimplemented!("{:?}", expression), + } + } + + fn write_block( + &mut self, + statements: &[crate::Statement], + ir_module: &crate::Module, + ir_function: &crate::Function, + function: &mut Function, + ) -> spirv::Word { + let mut block = Block::new(); + let id = self.generate_id(); + block.label = Some(super::instructions::instruction_label(id)); + + for statement in statements { + match *statement { + crate::Statement::Block(ref ir_block) => { + if !ir_block.is_empty() { + //TODO: link the block with `OpBranch` + self.write_block(ir_block, ir_module, ir_function, function); + } + } + crate::Statement::Return { value } => { + block.termination = Some(match ir_function.return_type { + Some(_) => { + let expression = &ir_function.expressions[value.unwrap()]; + let (id, lookup_ty) = self + .write_expression( + ir_module, + ir_function, + expression, + &mut block, + function, + ) + .unwrap(); + + let id = match *expression { + crate::Expression::LocalVariable(_) + | crate::Expression::GlobalVariable(_) => { + let load_id = self.generate_id(); + let value_ty_id = self.get_type_id(&ir_module.types, lookup_ty); + block.body.push(super::instructions::instruction_load( + value_ty_id, + load_id, + id, + None, + )); + load_id + } + + _ => id, + }; + super::instructions::instruction_return_value(id) + } + None => super::instructions::instruction_return(), + }); + } + crate::Statement::Store { pointer, value } => { + let pointer_expression = &ir_function.expressions[pointer]; + let value_expression = &ir_function.expressions[value]; + let (pointer_id, _) = self + .write_expression( + ir_module, + ir_function, + pointer_expression, + &mut block, + function, + ) + .unwrap(); + let (value_id, value_lookup_ty) = self + .write_expression( + ir_module, + ir_function, + value_expression, + &mut block, + function, + ) + .unwrap(); + + let value_id = match value_expression { + crate::Expression::LocalVariable(_) + | crate::Expression::GlobalVariable(_) => { + let load_id = self.generate_id(); + let value_ty_id = self.get_type_id(&ir_module.types, value_lookup_ty); + block.body.push(super::instructions::instruction_load( + value_ty_id, + load_id, + value_id, + None, + )); + load_id + } + _ => value_id, + }; + + block.body.push(super::instructions::instruction_store( + pointer_id, value_id, None, + )); + } + _ => unimplemented!("{:?}", statement), + } + } + + function.blocks.push(block); + id + } + + fn write_physical_layout(&mut self) { + self.physical_layout.bound = self.id_count + 1; + } + + fn write_logical_layout(&mut self, ir_module: &crate::Module) { + let id = self.generate_id(); + super::instructions::instruction_ext_inst_import(id, "GLSL.std.450") + .to_words(&mut self.logical_layout.ext_inst_imports); + + if self.writer_flags.contains(WriterFlags::DEBUG) { + self.debugs.push(super::instructions::instruction_source( + spirv::SourceLanguage::GLSL, + 450, + )); + } + + for (handle, ir_function) in ir_module.functions.iter() { + let id = self.write_function(ir_function, ir_module); + self.lookup_function.insert(handle, id); + } + + for (&(stage, ref name), ir_ep) in ir_module.entry_points.iter() { + let entry_point_instruction = self.write_entry_point(ir_ep, stage, name, ir_module); + entry_point_instruction.to_words(&mut self.logical_layout.entry_points); + } + + for capability in self.capabilities.iter() { + super::instructions::instruction_capability(*capability) + .to_words(&mut self.logical_layout.capabilities); + } + + let addressing_model = spirv::AddressingModel::Logical; + let memory_model = spirv::MemoryModel::GLSL450; + self.try_add_capabilities(addressing_model.required_capabilities()); + self.try_add_capabilities(memory_model.required_capabilities()); + + super::instructions::instruction_memory_model(addressing_model, memory_model) + .to_words(&mut self.logical_layout.memory_model); + + if self.writer_flags.contains(WriterFlags::DEBUG) { + for debug in self.debugs.iter() { + debug.to_words(&mut self.logical_layout.debugs); + } + } + + for annotation in self.annotations.iter() { + annotation.to_words(&mut self.logical_layout.annotations); + } + } + + pub fn write(&mut self, ir_module: &crate::Module) -> Vec<Word> { + let mut words: Vec<Word> = vec![]; + + self.write_logical_layout(ir_module); + self.write_physical_layout(); + + self.physical_layout.in_words(&mut words); + self.logical_layout.in_words(&mut words); + words + } +} + +#[cfg(test)] +mod tests { + use crate::back::spv::{Writer, WriterFlags}; + use crate::Header; + + #[test] + fn test_writer_generate_id() { + let mut writer = create_writer(); + + assert_eq!(writer.id_count, 0); + writer.generate_id(); + assert_eq!(writer.id_count, 1); + } + + #[test] + fn test_try_add_capabilities() { + let mut writer = create_writer(); + + assert_eq!(writer.capabilities.len(), 0); + writer.try_add_capabilities(&[spirv::Capability::Shader]); + assert_eq!(writer.capabilities.len(), 1); + + writer.try_add_capabilities(&[spirv::Capability::Shader]); + assert_eq!(writer.capabilities.len(), 1); + } + + #[test] + fn test_write_physical_layout() { + let mut writer = create_writer(); + assert_eq!(writer.physical_layout.bound, 0); + writer.write_physical_layout(); + assert_eq!(writer.physical_layout.bound, 1); + } + + fn create_writer() -> Writer { + let header = Header { + generator: 0, + version: (1, 0, 0), + }; + Writer::new(&header, WriterFlags::NONE) + } +} |