Source code
Revision control
Copy as Markdown
Other Tools
use alloc::{
format,
string::{String, ToString},
vec,
vec::Vec,
};
use core::fmt::Write;
use super::Error;
use super::ToWgslIfImplemented as _;
use crate::{back::wgsl::polyfill::InversePolyfill, common::wgsl::TypeContext};
use crate::{
back::{self, Baked},
common::{
self,
wgsl::{address_space_str, ToWgsl, TryToWgsl},
},
proc::{self, NameKey},
valid, Handle, Module, ShaderStage, TypeInner,
};
/// Shorthand result used internally by the backend
type BackendResult = Result<(), Error>;
enum Attribute {
Binding(u32),
BuiltIn(crate::BuiltIn),
Group(u32),
Invariant,
Interpolate(Option<crate::Interpolation>, Option<crate::Sampling>),
Location(u32),
BlendSrc(u32),
Stage(ShaderStage),
WorkGroupSize([u32; 3]),
}
/// The WGSL form that `write_expr_with_indirection` should use to render a Naga
/// expression.
///
/// Sometimes a Naga `Expression` alone doesn't provide enough information to
/// choose the right rendering for it in WGSL. For example, one natural WGSL
/// rendering of a Naga `LocalVariable(x)` expression might be `&x`, since
/// `LocalVariable` produces a pointer to the local variable's storage. But when
/// rendering a `Store` statement, the `pointer` operand must be the left hand
/// side of a WGSL assignment, so the proper rendering is `x`.
///
/// The caller of `write_expr_with_indirection` must provide an `Expected` value
/// to indicate how ambiguous expressions should be rendered.
#[derive(Clone, Copy, Debug)]
enum Indirection {
/// Render pointer-construction expressions as WGSL `ptr`-typed expressions.
///
/// This is the right choice for most cases. Whenever a Naga pointer
/// expression is not the `pointer` operand of a `Load` or `Store`, it
/// must be a WGSL pointer expression.
Ordinary,
/// Render pointer-construction expressions as WGSL reference-typed
/// expressions.
///
/// For example, this is the right choice for the `pointer` operand when
/// rendering a `Store` statement as a WGSL assignment.
Reference,
}
bitflags::bitflags! {
#[cfg_attr(feature = "serialize", derive(serde::Serialize))]
#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
#[derive(Clone, Copy, Debug, Eq, PartialEq)]
pub struct WriterFlags: u32 {
/// Always annotate the type information instead of inferring.
const EXPLICIT_TYPES = 0x1;
}
}
pub struct Writer<W> {
out: W,
flags: WriterFlags,
names: crate::FastHashMap<NameKey, String>,
namer: proc::Namer,
named_expressions: crate::NamedExpressions,
required_polyfills: crate::FastIndexSet<InversePolyfill>,
}
impl<W: Write> Writer<W> {
pub fn new(out: W, flags: WriterFlags) -> Self {
Writer {
out,
flags,
names: crate::FastHashMap::default(),
namer: proc::Namer::default(),
named_expressions: crate::NamedExpressions::default(),
required_polyfills: crate::FastIndexSet::default(),
}
}
fn reset(&mut self, module: &Module) {
self.names.clear();
self.namer.reset(
module,
&crate::keywords::wgsl::RESERVED_SET,
// an identifier must not start with two underscore
&[],
&["__", "_naga"],
&mut self.names,
);
self.named_expressions.clear();
self.required_polyfills.clear();
}
/// Determine if `ty` is the Naga IR presentation of a WGSL builtin type.
///
/// Return true if `ty` refers to the Naga IR form of a WGSL builtin type
/// like `__atomic_compare_exchange_result`.
///
/// Even though the module may use the type, the WGSL backend should avoid
/// emitting a definition for it, since it is [predeclared] in WGSL.
///
/// This also covers types like [`NagaExternalTextureParams`], which other
/// backends use to lower WGSL constructs like external textures to their
/// implementations. WGSL can express these directly, so the types need not
/// be emitted.
///
/// [`NagaExternalTextureParams`]: crate::ir::SpecialTypes::external_texture_params
fn is_builtin_wgsl_struct(&self, module: &Module, ty: Handle<crate::Type>) -> bool {
module
.special_types
.predeclared_types
.values()
.any(|t| *t == ty)
|| Some(ty) == module.special_types.external_texture_params
|| Some(ty) == module.special_types.external_texture_transfer_function
}
pub fn write(&mut self, module: &Module, info: &valid::ModuleInfo) -> BackendResult {
if !module.overrides.is_empty() {
return Err(Error::Unimplemented(
"Pipeline constants are not yet supported for this back-end".to_string(),
));
}
self.reset(module);
// Write all `enable` declarations
self.write_enable_declarations(module)?;
// Write all structs
for (handle, ty) in module.types.iter() {
if let TypeInner::Struct { ref members, .. } = ty.inner {
{
if !self.is_builtin_wgsl_struct(module, handle) {
self.write_struct(module, handle, members)?;
writeln!(self.out)?;
}
}
}
}
// Write all named constants
let mut constants = module
.constants
.iter()
.filter(|&(_, c)| c.name.is_some())
.peekable();
while let Some((handle, _)) = constants.next() {
self.write_global_constant(module, handle)?;
// Add extra newline for readability on last iteration
if constants.peek().is_none() {
writeln!(self.out)?;
}
}
// Write all globals
for (ty, global) in module.global_variables.iter() {
self.write_global(module, global, ty)?;
}
if !module.global_variables.is_empty() {
// Add extra newline for readability
writeln!(self.out)?;
}
// Write all regular functions
for (handle, function) in module.functions.iter() {
let fun_info = &info[handle];
let func_ctx = back::FunctionCtx {
ty: back::FunctionType::Function(handle),
info: fun_info,
expressions: &function.expressions,
named_expressions: &function.named_expressions,
};
// Write the function
self.write_function(module, function, &func_ctx)?;
writeln!(self.out)?;
}
// Write all entry points
for (index, ep) in module.entry_points.iter().enumerate() {
let attributes = match ep.stage {
ShaderStage::Vertex | ShaderStage::Fragment => vec![Attribute::Stage(ep.stage)],
ShaderStage::Compute => vec![
Attribute::Stage(ShaderStage::Compute),
Attribute::WorkGroupSize(ep.workgroup_size),
],
ShaderStage::Task | ShaderStage::Mesh => unreachable!(),
};
self.write_attributes(&attributes)?;
// Add a newline after attribute
writeln!(self.out)?;
let func_ctx = back::FunctionCtx {
ty: back::FunctionType::EntryPoint(index as u16),
info: info.get_entry_point(index),
expressions: &ep.function.expressions,
named_expressions: &ep.function.named_expressions,
};
self.write_function(module, &ep.function, &func_ctx)?;
if index < module.entry_points.len() - 1 {
writeln!(self.out)?;
}
}
// Write any polyfills that were required.
for polyfill in &self.required_polyfills {
writeln!(self.out)?;
write!(self.out, "{}", polyfill.source)?;
writeln!(self.out)?;
}
Ok(())
}
/// Helper method which writes all the `enable` declarations
/// needed for a module.
fn write_enable_declarations(&mut self, module: &Module) -> BackendResult {
let mut needs_f16 = false;
let mut needs_dual_source_blending = false;
let mut needs_clip_distances = false;
// Determine which `enable` declarations are needed
for (_, ty) in module.types.iter() {
match ty.inner {
TypeInner::Scalar(scalar)
| TypeInner::Vector { scalar, .. }
| TypeInner::Matrix { scalar, .. } => {
needs_f16 |= scalar == crate::Scalar::F16;
}
TypeInner::Struct { ref members, .. } => {
for binding in members.iter().filter_map(|m| m.binding.as_ref()) {
match *binding {
crate::Binding::Location {
blend_src: Some(_), ..
} => {
needs_dual_source_blending = true;
}
crate::Binding::BuiltIn(crate::BuiltIn::ClipDistance) => {
needs_clip_distances = true;
}
_ => {}
}
}
}
_ => {}
}
}
// Write required declarations
let mut any_written = false;
if needs_f16 {
writeln!(self.out, "enable f16;")?;
any_written = true;
}
if needs_dual_source_blending {
writeln!(self.out, "enable dual_source_blending;")?;
any_written = true;
}
if needs_clip_distances {
writeln!(self.out, "enable clip_distances;")?;
any_written = true;
}
if any_written {
// Empty line for readability
writeln!(self.out)?;
}
Ok(())
}
/// Helper method used to write
///
/// # Notes
/// Ends in a newline
fn write_function(
&mut self,
module: &Module,
func: &crate::Function,
func_ctx: &back::FunctionCtx<'_>,
) -> BackendResult {
let func_name = match func_ctx.ty {
back::FunctionType::EntryPoint(index) => &self.names[&NameKey::EntryPoint(index)],
back::FunctionType::Function(handle) => &self.names[&NameKey::Function(handle)],
};
// Write function name
write!(self.out, "fn {func_name}(")?;
// Write function arguments
for (index, arg) in func.arguments.iter().enumerate() {
// Write argument attribute if a binding is present
if let Some(ref binding) = arg.binding {
self.write_attributes(&map_binding_to_attribute(binding))?;
}
// Write argument name
let argument_name = &self.names[&func_ctx.argument_key(index as u32)];
write!(self.out, "{argument_name}: ")?;
// Write argument type
self.write_type(module, arg.ty)?;
if index < func.arguments.len() - 1 {
// Add a separator between args
write!(self.out, ", ")?;
}
}
write!(self.out, ")")?;
// Write function return type
if let Some(ref result) = func.result {
write!(self.out, " -> ")?;
if let Some(ref binding) = result.binding {
self.write_attributes(&map_binding_to_attribute(binding))?;
}
self.write_type(module, result.ty)?;
}
write!(self.out, " {{")?;
writeln!(self.out)?;
// Write function local variables
for (handle, local) in func.local_variables.iter() {
// Write indentation (only for readability)
write!(self.out, "{}", back::INDENT)?;
// Write the local name
// The leading space is important
write!(self.out, "var {}: ", self.names[&func_ctx.name_key(handle)])?;
// Write the local type
self.write_type(module, local.ty)?;
// Write the local initializer if needed
if let Some(init) = local.init {
// Put the equal signal only if there's a initializer
// The leading and trailing spaces aren't needed but help with readability
write!(self.out, " = ")?;
// Write the constant
// `write_constant` adds no trailing or leading space/newline
self.write_expr(module, init, func_ctx)?;
}
// Finish the local with `;` and add a newline (only for readability)
writeln!(self.out, ";")?
}
if !func.local_variables.is_empty() {
writeln!(self.out)?;
}
// Write the function body (statement list)
for sta in func.body.iter() {
// The indentation should always be 1 when writing the function body
self.write_stmt(module, sta, func_ctx, back::Level(1))?;
}
writeln!(self.out, "}}")?;
self.named_expressions.clear();
Ok(())
}
/// Helper method to write a attribute
fn write_attributes(&mut self, attributes: &[Attribute]) -> BackendResult {
for attribute in attributes {
match *attribute {
Attribute::Location(id) => write!(self.out, "@location({id}) ")?,
Attribute::BlendSrc(blend_src) => write!(self.out, "@blend_src({blend_src}) ")?,
Attribute::BuiltIn(builtin_attrib) => {
let builtin = builtin_attrib.to_wgsl_if_implemented()?;
write!(self.out, "@builtin({builtin}) ")?;
}
Attribute::Stage(shader_stage) => {
let stage_str = match shader_stage {
ShaderStage::Vertex => "vertex",
ShaderStage::Fragment => "fragment",
ShaderStage::Compute => "compute",
ShaderStage::Task | ShaderStage::Mesh => unreachable!(),
};
write!(self.out, "@{stage_str} ")?;
}
Attribute::WorkGroupSize(size) => {
write!(
self.out,
"@workgroup_size({}, {}, {}) ",
size[0], size[1], size[2]
)?;
}
Attribute::Binding(id) => write!(self.out, "@binding({id}) ")?,
Attribute::Group(id) => write!(self.out, "@group({id}) ")?,
Attribute::Invariant => write!(self.out, "@invariant ")?,
Attribute::Interpolate(interpolation, sampling) => {
if sampling.is_some() && sampling != Some(crate::Sampling::Center) {
let interpolation = interpolation
.unwrap_or(crate::Interpolation::Perspective)
.to_wgsl();
let sampling = sampling.unwrap_or(crate::Sampling::Center).to_wgsl();
write!(self.out, "@interpolate({interpolation}, {sampling}) ")?;
} else if interpolation.is_some()
&& interpolation != Some(crate::Interpolation::Perspective)
{
let interpolation = interpolation
.unwrap_or(crate::Interpolation::Perspective)
.to_wgsl();
write!(self.out, "@interpolate({interpolation}) ")?;
}
}
};
}
Ok(())
}
/// Helper method used to write structs
/// Write the full declaration of a struct type.
///
/// Write out a definition of the struct type referred to by
/// `handle` in `module`. The output will be an instance of the
/// `struct_decl` production in the WGSL grammar.
///
/// Use `members` as the list of `handle`'s members. (This
/// function is usually called after matching a `TypeInner`, so
/// the callers already have the members at hand.)
fn write_struct(
&mut self,
module: &Module,
handle: Handle<crate::Type>,
members: &[crate::StructMember],
) -> BackendResult {
write!(self.out, "struct {}", self.names[&NameKey::Type(handle)])?;
write!(self.out, " {{")?;
writeln!(self.out)?;
for (index, member) in members.iter().enumerate() {
// The indentation is only for readability
write!(self.out, "{}", back::INDENT)?;
if let Some(ref binding) = member.binding {
self.write_attributes(&map_binding_to_attribute(binding))?;
}
// Write struct member name and type
let member_name = &self.names[&NameKey::StructMember(handle, index as u32)];
write!(self.out, "{member_name}: ")?;
self.write_type(module, member.ty)?;
write!(self.out, ",")?;
writeln!(self.out)?;
}
writeln!(self.out, "}}")?;
Ok(())
}
fn write_type(&mut self, module: &Module, ty: Handle<crate::Type>) -> BackendResult {
// This actually can't be factored out into a nice constructor method,
// because the borrow checker needs to be able to see that the borrows
// of `self.names` and `self.out` are disjoint.
let type_context = WriterTypeContext {
module,
names: &self.names,
};
type_context.write_type(ty, &mut self.out)?;
Ok(())
}
fn write_type_resolution(
&mut self,
module: &Module,
resolution: &proc::TypeResolution,
) -> BackendResult {
// This actually can't be factored out into a nice constructor method,
// because the borrow checker needs to be able to see that the borrows
// of `self.names` and `self.out` are disjoint.
let type_context = WriterTypeContext {
module,
names: &self.names,
};
type_context.write_type_resolution(resolution, &mut self.out)?;
Ok(())
}
/// Helper method used to write statements
///
/// # Notes
/// Always adds a newline
fn write_stmt(
&mut self,
module: &Module,
stmt: &crate::Statement,
func_ctx: &back::FunctionCtx<'_>,
level: back::Level,
) -> BackendResult {
use crate::{Expression, Statement};
match *stmt {
Statement::Emit(ref range) => {
for handle in range.clone() {
let info = &func_ctx.info[handle];
let expr_name = if let Some(name) = func_ctx.named_expressions.get(&handle) {
// Front end provides names for all variables at the start of writing.
// But we write them to step by step. We need to recache them
// Otherwise, we could accidentally write variable name instead of full expression.
// Also, we use sanitized names! It defense backend from generating variable with name from reserved keywords.
Some(self.namer.call(name))
} else {
let expr = &func_ctx.expressions[handle];
let min_ref_count = expr.bake_ref_count();
// Forcefully creating baking expressions in some cases to help with readability
let required_baking_expr = match *expr {
Expression::ImageLoad { .. }
| Expression::ImageQuery { .. }
| Expression::ImageSample { .. } => true,
_ => false,
};
if min_ref_count <= info.ref_count || required_baking_expr {
Some(Baked(handle).to_string())
} else {
None
}
};
if let Some(name) = expr_name {
write!(self.out, "{level}")?;
self.start_named_expr(module, handle, func_ctx, &name)?;
self.write_expr(module, handle, func_ctx)?;
self.named_expressions.insert(handle, name);
writeln!(self.out, ";")?;
}
}
}
// TODO: copy-paste from glsl-out
Statement::If {
condition,
ref accept,
ref reject,
} => {
write!(self.out, "{level}")?;
write!(self.out, "if ")?;
self.write_expr(module, condition, func_ctx)?;
writeln!(self.out, " {{")?;
let l2 = level.next();
for sta in accept {
// Increase indentation to help with readability
self.write_stmt(module, sta, func_ctx, l2)?;
}
// If there are no statements in the reject block we skip writing it
// This is only for readability
if !reject.is_empty() {
writeln!(self.out, "{level}}} else {{")?;
for sta in reject {
// Increase indentation to help with readability
self.write_stmt(module, sta, func_ctx, l2)?;
}
}
writeln!(self.out, "{level}}}")?
}
Statement::Return { value } => {
write!(self.out, "{level}")?;
write!(self.out, "return")?;
if let Some(return_value) = value {
// The leading space is important
write!(self.out, " ")?;
self.write_expr(module, return_value, func_ctx)?;
}
writeln!(self.out, ";")?;
}
// TODO: copy-paste from glsl-out
Statement::Kill => {
write!(self.out, "{level}")?;
writeln!(self.out, "discard;")?
}
Statement::Store { pointer, value } => {
write!(self.out, "{level}")?;
let is_atomic_pointer = func_ctx
.resolve_type(pointer, &module.types)
.is_atomic_pointer(&module.types);
if is_atomic_pointer {
write!(self.out, "atomicStore(")?;
self.write_expr(module, pointer, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, value, func_ctx)?;
write!(self.out, ")")?;
} else {
self.write_expr_with_indirection(
module,
pointer,
func_ctx,
Indirection::Reference,
)?;
write!(self.out, " = ")?;
self.write_expr(module, value, func_ctx)?;
}
writeln!(self.out, ";")?
}
Statement::Call {
function,
ref arguments,
result,
} => {
write!(self.out, "{level}")?;
if let Some(expr) = result {
let name = Baked(expr).to_string();
self.start_named_expr(module, expr, func_ctx, &name)?;
self.named_expressions.insert(expr, name);
}
let func_name = &self.names[&NameKey::Function(function)];
write!(self.out, "{func_name}(")?;
for (index, &argument) in arguments.iter().enumerate() {
if index != 0 {
write!(self.out, ", ")?;
}
self.write_expr(module, argument, func_ctx)?;
}
writeln!(self.out, ");")?
}
Statement::Atomic {
pointer,
ref fun,
value,
result,
} => {
write!(self.out, "{level}")?;
if let Some(result) = result {
let res_name = Baked(result).to_string();
self.start_named_expr(module, result, func_ctx, &res_name)?;
self.named_expressions.insert(result, res_name);
}
let fun_str = fun.to_wgsl();
write!(self.out, "atomic{fun_str}(")?;
self.write_expr(module, pointer, func_ctx)?;
if let crate::AtomicFunction::Exchange { compare: Some(cmp) } = *fun {
write!(self.out, ", ")?;
self.write_expr(module, cmp, func_ctx)?;
}
write!(self.out, ", ")?;
self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ");")?
}
Statement::ImageAtomic {
image,
coordinate,
array_index,
ref fun,
value,
} => {
write!(self.out, "{level}")?;
let fun_str = fun.to_wgsl();
write!(self.out, "textureAtomic{fun_str}(")?;
self.write_expr(module, image, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, coordinate, func_ctx)?;
if let Some(array_index_expr) = array_index {
write!(self.out, ", ")?;
self.write_expr(module, array_index_expr, func_ctx)?;
}
write!(self.out, ", ")?;
self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ");")?;
}
Statement::WorkGroupUniformLoad { pointer, result } => {
write!(self.out, "{level}")?;
// TODO: Obey named expressions here.
let res_name = Baked(result).to_string();
self.start_named_expr(module, result, func_ctx, &res_name)?;
self.named_expressions.insert(result, res_name);
write!(self.out, "workgroupUniformLoad(")?;
self.write_expr(module, pointer, func_ctx)?;
writeln!(self.out, ");")?;
}
Statement::ImageStore {
image,
coordinate,
array_index,
value,
} => {
write!(self.out, "{level}")?;
write!(self.out, "textureStore(")?;
self.write_expr(module, image, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, coordinate, func_ctx)?;
if let Some(array_index_expr) = array_index {
write!(self.out, ", ")?;
self.write_expr(module, array_index_expr, func_ctx)?;
}
write!(self.out, ", ")?;
self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ");")?;
}
// TODO: copy-paste from glsl-out
Statement::Block(ref block) => {
write!(self.out, "{level}")?;
writeln!(self.out, "{{")?;
for sta in block.iter() {
// Increase the indentation to help with readability
self.write_stmt(module, sta, func_ctx, level.next())?
}
writeln!(self.out, "{level}}}")?
}
Statement::Switch {
selector,
ref cases,
} => {
// Start the switch
write!(self.out, "{level}")?;
write!(self.out, "switch ")?;
self.write_expr(module, selector, func_ctx)?;
writeln!(self.out, " {{")?;
let l2 = level.next();
let mut new_case = true;
for case in cases {
if case.fall_through && !case.body.is_empty() {
// TODO: we could do the same workaround as we did for the HLSL backend
return Err(Error::Unimplemented(
"fall-through switch case block".into(),
));
}
match case.value {
crate::SwitchValue::I32(value) => {
if new_case {
write!(self.out, "{l2}case ")?;
}
write!(self.out, "{value}")?;
}
crate::SwitchValue::U32(value) => {
if new_case {
write!(self.out, "{l2}case ")?;
}
write!(self.out, "{value}u")?;
}
crate::SwitchValue::Default => {
if new_case {
if case.fall_through {
write!(self.out, "{l2}case ")?;
} else {
write!(self.out, "{l2}")?;
}
}
write!(self.out, "default")?;
}
}
new_case = !case.fall_through;
if case.fall_through {
write!(self.out, ", ")?;
} else {
writeln!(self.out, ": {{")?;
}
for sta in case.body.iter() {
self.write_stmt(module, sta, func_ctx, l2.next())?;
}
if !case.fall_through {
writeln!(self.out, "{l2}}}")?;
}
}
writeln!(self.out, "{level}}}")?
}
Statement::Loop {
ref body,
ref continuing,
break_if,
} => {
write!(self.out, "{level}")?;
writeln!(self.out, "loop {{")?;
let l2 = level.next();
for sta in body.iter() {
self.write_stmt(module, sta, func_ctx, l2)?;
}
// The continuing is optional so we don't need to write it if
// it is empty, but the `break if` counts as a continuing statement
// so even if `continuing` is empty we must generate it if a
// `break if` exists
if !continuing.is_empty() || break_if.is_some() {
writeln!(self.out, "{l2}continuing {{")?;
for sta in continuing.iter() {
self.write_stmt(module, sta, func_ctx, l2.next())?;
}
// The `break if` is always the last
// statement of the `continuing` block
if let Some(condition) = break_if {
// The trailing space is important
write!(self.out, "{}break if ", l2.next())?;
self.write_expr(module, condition, func_ctx)?;
// Close the `break if` statement
writeln!(self.out, ";")?;
}
writeln!(self.out, "{l2}}}")?;
}
writeln!(self.out, "{level}}}")?
}
Statement::Break => {
writeln!(self.out, "{level}break;")?;
}
Statement::Continue => {
writeln!(self.out, "{level}continue;")?;
}
Statement::ControlBarrier(barrier) | Statement::MemoryBarrier(barrier) => {
if barrier.contains(crate::Barrier::STORAGE) {
writeln!(self.out, "{level}storageBarrier();")?;
}
if barrier.contains(crate::Barrier::WORK_GROUP) {
writeln!(self.out, "{level}workgroupBarrier();")?;
}
if barrier.contains(crate::Barrier::SUB_GROUP) {
writeln!(self.out, "{level}subgroupBarrier();")?;
}
if barrier.contains(crate::Barrier::TEXTURE) {
writeln!(self.out, "{level}textureBarrier();")?;
}
}
Statement::RayQuery { .. } => unreachable!(),
Statement::SubgroupBallot { result, predicate } => {
write!(self.out, "{level}")?;
let res_name = Baked(result).to_string();
self.start_named_expr(module, result, func_ctx, &res_name)?;
self.named_expressions.insert(result, res_name);
write!(self.out, "subgroupBallot(")?;
if let Some(predicate) = predicate {
self.write_expr(module, predicate, func_ctx)?;
}
writeln!(self.out, ");")?;
}
Statement::SubgroupCollectiveOperation {
op,
collective_op,
argument,
result,
} => {
write!(self.out, "{level}")?;
let res_name = Baked(result).to_string();
self.start_named_expr(module, result, func_ctx, &res_name)?;
self.named_expressions.insert(result, res_name);
match (collective_op, op) {
(crate::CollectiveOperation::Reduce, crate::SubgroupOperation::All) => {
write!(self.out, "subgroupAll(")?
}
(crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Any) => {
write!(self.out, "subgroupAny(")?
}
(crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Add) => {
write!(self.out, "subgroupAdd(")?
}
(crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Mul) => {
write!(self.out, "subgroupMul(")?
}
(crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Max) => {
write!(self.out, "subgroupMax(")?
}
(crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Min) => {
write!(self.out, "subgroupMin(")?
}
(crate::CollectiveOperation::Reduce, crate::SubgroupOperation::And) => {
write!(self.out, "subgroupAnd(")?
}
(crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Or) => {
write!(self.out, "subgroupOr(")?
}
(crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Xor) => {
write!(self.out, "subgroupXor(")?
}
(crate::CollectiveOperation::ExclusiveScan, crate::SubgroupOperation::Add) => {
write!(self.out, "subgroupExclusiveAdd(")?
}
(crate::CollectiveOperation::ExclusiveScan, crate::SubgroupOperation::Mul) => {
write!(self.out, "subgroupExclusiveMul(")?
}
(crate::CollectiveOperation::InclusiveScan, crate::SubgroupOperation::Add) => {
write!(self.out, "subgroupInclusiveAdd(")?
}
(crate::CollectiveOperation::InclusiveScan, crate::SubgroupOperation::Mul) => {
write!(self.out, "subgroupInclusiveMul(")?
}
_ => unimplemented!(),
}
self.write_expr(module, argument, func_ctx)?;
writeln!(self.out, ");")?;
}
Statement::SubgroupGather {
mode,
argument,
result,
} => {
write!(self.out, "{level}")?;
let res_name = Baked(result).to_string();
self.start_named_expr(module, result, func_ctx, &res_name)?;
self.named_expressions.insert(result, res_name);
match mode {
crate::GatherMode::BroadcastFirst => {
write!(self.out, "subgroupBroadcastFirst(")?;
}
crate::GatherMode::Broadcast(_) => {
write!(self.out, "subgroupBroadcast(")?;
}
crate::GatherMode::Shuffle(_) => {
write!(self.out, "subgroupShuffle(")?;
}
crate::GatherMode::ShuffleDown(_) => {
write!(self.out, "subgroupShuffleDown(")?;
}
crate::GatherMode::ShuffleUp(_) => {
write!(self.out, "subgroupShuffleUp(")?;
}
crate::GatherMode::ShuffleXor(_) => {
write!(self.out, "subgroupShuffleXor(")?;
}
crate::GatherMode::QuadBroadcast(_) => {
write!(self.out, "quadBroadcast(")?;
}
crate::GatherMode::QuadSwap(direction) => match direction {
crate::Direction::X => {
write!(self.out, "quadSwapX(")?;
}
crate::Direction::Y => {
write!(self.out, "quadSwapY(")?;
}
crate::Direction::Diagonal => {
write!(self.out, "quadSwapDiagonal(")?;
}
},
}
self.write_expr(module, argument, func_ctx)?;
match mode {
crate::GatherMode::BroadcastFirst => {}
crate::GatherMode::Broadcast(index)
| crate::GatherMode::Shuffle(index)
| crate::GatherMode::ShuffleDown(index)
| crate::GatherMode::ShuffleUp(index)
| crate::GatherMode::ShuffleXor(index)
| crate::GatherMode::QuadBroadcast(index) => {
write!(self.out, ", ")?;
self.write_expr(module, index, func_ctx)?;
}
crate::GatherMode::QuadSwap(_) => {}
}
writeln!(self.out, ");")?;
}
}
Ok(())
}
/// Return the sort of indirection that `expr`'s plain form evaluates to.
///
/// An expression's 'plain form' is the most general rendition of that
/// expression into WGSL, lacking `&` or `*` operators:
///
/// - The plain form of `LocalVariable(x)` is simply `x`, which is a reference
/// to the local variable's storage.
///
/// - The plain form of `GlobalVariable(g)` is simply `g`, which is usually a
/// reference to the global variable's storage. However, globals in the
/// `Handle` address space are immutable, and `GlobalVariable` expressions for
/// those produce the value directly, not a pointer to it. Such
/// `GlobalVariable` expressions are `Ordinary`.
///
/// - `Access` and `AccessIndex` are `Reference` when their `base` operand is a
/// pointer. If they are applied directly to a composite value, they are
/// `Ordinary`.
///
/// Note that `FunctionArgument` expressions are never `Reference`, even when
/// the argument's type is `Pointer`. `FunctionArgument` always evaluates to the
/// argument's value directly, so any pointer it produces is merely the value
/// passed by the caller.
fn plain_form_indirection(
&self,
expr: Handle<crate::Expression>,
module: &Module,
func_ctx: &back::FunctionCtx<'_>,
) -> Indirection {
use crate::Expression as Ex;
// Named expressions are `let` expressions, which apply the Load Rule,
// so if their type is a Naga pointer, then that must be a WGSL pointer
// as well.
if self.named_expressions.contains_key(&expr) {
return Indirection::Ordinary;
}
match func_ctx.expressions[expr] {
Ex::LocalVariable(_) => Indirection::Reference,
Ex::GlobalVariable(handle) => {
let global = &module.global_variables[handle];
match global.space {
crate::AddressSpace::Handle => Indirection::Ordinary,
_ => Indirection::Reference,
}
}
Ex::Access { base, .. } | Ex::AccessIndex { base, .. } => {
let base_ty = func_ctx.resolve_type(base, &module.types);
match *base_ty {
TypeInner::Pointer { .. } | TypeInner::ValuePointer { .. } => {
Indirection::Reference
}
_ => Indirection::Ordinary,
}
}
_ => Indirection::Ordinary,
}
}
fn start_named_expr(
&mut self,
module: &Module,
handle: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx,
name: &str,
) -> BackendResult {
// Write variable name
write!(self.out, "let {name}")?;
if self.flags.contains(WriterFlags::EXPLICIT_TYPES) {
write!(self.out, ": ")?;
// Write variable type
self.write_type_resolution(module, &func_ctx.info[handle].ty)?;
}
write!(self.out, " = ")?;
Ok(())
}
/// Write the ordinary WGSL form of `expr`.
///
/// See `write_expr_with_indirection` for details.
fn write_expr(
&mut self,
module: &Module,
expr: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
) -> BackendResult {
self.write_expr_with_indirection(module, expr, func_ctx, Indirection::Ordinary)
}
/// Write `expr` as a WGSL expression with the requested indirection.
///
/// In terms of the WGSL grammar, the resulting expression is a
/// `singular_expression`. It may be parenthesized. This makes it suitable
/// for use as the operand of a unary or binary operator without worrying
/// about precedence.
///
/// This does not produce newlines or indentation.
///
/// The `requested` argument indicates (roughly) whether Naga
/// `Pointer`-valued expressions represent WGSL references or pointers. See
/// `Indirection` for details.
fn write_expr_with_indirection(
&mut self,
module: &Module,
expr: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
requested: Indirection,
) -> BackendResult {
// If the plain form of the expression is not what we need, emit the
// operator necessary to correct that.
let plain = self.plain_form_indirection(expr, module, func_ctx);
match (requested, plain) {
(Indirection::Ordinary, Indirection::Reference) => {
write!(self.out, "(&")?;
self.write_expr_plain_form(module, expr, func_ctx, plain)?;
write!(self.out, ")")?;
}
(Indirection::Reference, Indirection::Ordinary) => {
write!(self.out, "(*")?;
self.write_expr_plain_form(module, expr, func_ctx, plain)?;
write!(self.out, ")")?;
}
(_, _) => self.write_expr_plain_form(module, expr, func_ctx, plain)?,
}
Ok(())
}
fn write_const_expression(
&mut self,
module: &Module,
expr: Handle<crate::Expression>,
arena: &crate::Arena<crate::Expression>,
) -> BackendResult {
self.write_possibly_const_expression(module, expr, arena, |writer, expr| {
writer.write_const_expression(module, expr, arena)
})
}
fn write_possibly_const_expression<E>(
&mut self,
module: &Module,
expr: Handle<crate::Expression>,
expressions: &crate::Arena<crate::Expression>,
write_expression: E,
) -> BackendResult
where
E: Fn(&mut Self, Handle<crate::Expression>) -> BackendResult,
{
use crate::Expression;
match expressions[expr] {
Expression::Literal(literal) => match literal {
crate::Literal::F16(value) => write!(self.out, "{value}h")?,
crate::Literal::F32(value) => write!(self.out, "{value}f")?,
crate::Literal::U32(value) => write!(self.out, "{value}u")?,
crate::Literal::I32(value) => {
// `-2147483648i` is not valid WGSL. The most negative `i32`
// value can only be expressed in WGSL using AbstractInt and
// a unary negation operator.
if value == i32::MIN {
write!(self.out, "i32({value})")?;
} else {
write!(self.out, "{value}i")?;
}
}
crate::Literal::Bool(value) => write!(self.out, "{value}")?,
crate::Literal::F64(value) => write!(self.out, "{value:?}lf")?,
crate::Literal::I64(value) => {
// `-9223372036854775808li` is not valid WGSL. Nor can we simply use the
// AbstractInt trick above, as AbstractInt also cannot represent
// `9223372036854775808`. Instead construct the second most negative
// AbstractInt, subtract one from it, then cast to i64.
if value == i64::MIN {
write!(self.out, "i64({} - 1)", value + 1)?;
} else {
write!(self.out, "{value}li")?;
}
}
crate::Literal::U64(value) => write!(self.out, "{value:?}lu")?,
crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => {
return Err(Error::Custom(
"Abstract types should not appear in IR presented to backends".into(),
));
}
},
Expression::Constant(handle) => {
let constant = &module.constants[handle];
if constant.name.is_some() {
write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
} else {
self.write_const_expression(module, constant.init, &module.global_expressions)?;
}
}
Expression::ZeroValue(ty) => {
self.write_type(module, ty)?;
write!(self.out, "()")?;
}
Expression::Compose { ty, ref components } => {
self.write_type(module, ty)?;
write!(self.out, "(")?;
for (index, component) in components.iter().enumerate() {
if index != 0 {
write!(self.out, ", ")?;
}
write_expression(self, *component)?;
}
write!(self.out, ")")?
}
Expression::Splat { size, value } => {
let size = common::vector_size_str(size);
write!(self.out, "vec{size}(")?;
write_expression(self, value)?;
write!(self.out, ")")?;
}
_ => unreachable!(),
}
Ok(())
}
/// Write the 'plain form' of `expr`.
///
/// An expression's 'plain form' is the most general rendition of that
/// expression into WGSL, lacking `&` or `*` operators. The plain forms of
/// `LocalVariable(x)` and `GlobalVariable(g)` are simply `x` and `g`. Such
/// Naga expressions represent both WGSL pointers and references; it's the
/// caller's responsibility to distinguish those cases appropriately.
fn write_expr_plain_form(
&mut self,
module: &Module,
expr: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
indirection: Indirection,
) -> BackendResult {
use crate::Expression;
if let Some(name) = self.named_expressions.get(&expr) {
write!(self.out, "{name}")?;
return Ok(());
}
let expression = &func_ctx.expressions[expr];
// Write the plain WGSL form of a Naga expression.
//
// The plain form of `LocalVariable` and `GlobalVariable` expressions is
// simply the variable name; `*` and `&` operators are never emitted.
//
// The plain form of `Access` and `AccessIndex` expressions are WGSL
// `postfix_expression` forms for member/component access and
// subscripting.
match *expression {
Expression::Literal(_)
| Expression::Constant(_)
| Expression::ZeroValue(_)
| Expression::Compose { .. }
| Expression::Splat { .. } => {
self.write_possibly_const_expression(
module,
expr,
func_ctx.expressions,
|writer, expr| writer.write_expr(module, expr, func_ctx),
)?;
}
Expression::Override(_) => unreachable!(),
Expression::FunctionArgument(pos) => {
let name_key = func_ctx.argument_key(pos);
let name = &self.names[&name_key];
write!(self.out, "{name}")?;
}
Expression::Binary { op, left, right } => {
write!(self.out, "(")?;
self.write_expr(module, left, func_ctx)?;
write!(self.out, " {} ", back::binary_operation_str(op))?;
self.write_expr(module, right, func_ctx)?;
write!(self.out, ")")?;
}
Expression::Access { base, index } => {
self.write_expr_with_indirection(module, base, func_ctx, indirection)?;
write!(self.out, "[")?;
self.write_expr(module, index, func_ctx)?;
write!(self.out, "]")?
}
Expression::AccessIndex { base, index } => {
let base_ty_res = &func_ctx.info[base].ty;
let mut resolved = base_ty_res.inner_with(&module.types);
self.write_expr_with_indirection(module, base, func_ctx, indirection)?;
let base_ty_handle = match *resolved {
TypeInner::Pointer { base, space: _ } => {
resolved = &module.types[base].inner;
Some(base)
}
_ => base_ty_res.handle(),
};
match *resolved {
TypeInner::Vector { .. } => {
// Write vector access as a swizzle
write!(self.out, ".{}", back::COMPONENTS[index as usize])?
}
TypeInner::Matrix { .. }
| TypeInner::Array { .. }
| TypeInner::BindingArray { .. }
| TypeInner::ValuePointer { .. } => write!(self.out, "[{index}]")?,
TypeInner::Struct { .. } => {
// This will never panic in case the type is a `Struct`, this is not true
// for other types so we can only check while inside this match arm
let ty = base_ty_handle.unwrap();
write!(
self.out,
".{}",
&self.names[&NameKey::StructMember(ty, index)]
)?
}
ref other => return Err(Error::Custom(format!("Cannot index {other:?}"))),
}
}
Expression::ImageSample {
image,
sampler,
gather: None,
coordinate,
array_index,
offset,
level,
depth_ref,
clamp_to_edge,
} => {
use crate::SampleLevel as Sl;
let suffix_cmp = match depth_ref {
Some(_) => "Compare",
None => "",
};
let suffix_level = match level {
Sl::Auto => "",
Sl::Zero if clamp_to_edge => "BaseClampToEdge",
Sl::Zero | Sl::Exact(_) => "Level",
Sl::Bias(_) => "Bias",
Sl::Gradient { .. } => "Grad",
};
write!(self.out, "textureSample{suffix_cmp}{suffix_level}(")?;
self.write_expr(module, image, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, sampler, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, coordinate, func_ctx)?;
if let Some(array_index) = array_index {
write!(self.out, ", ")?;
self.write_expr(module, array_index, func_ctx)?;
}
if let Some(depth_ref) = depth_ref {
write!(self.out, ", ")?;
self.write_expr(module, depth_ref, func_ctx)?;
}
match level {
Sl::Auto => {}
Sl::Zero => {
// Level 0 is implied for depth comparison and BaseClampToEdge
if depth_ref.is_none() && !clamp_to_edge {
write!(self.out, ", 0.0")?;
}
}
Sl::Exact(expr) => {
write!(self.out, ", ")?;
self.write_expr(module, expr, func_ctx)?;
}
Sl::Bias(expr) => {
write!(self.out, ", ")?;
self.write_expr(module, expr, func_ctx)?;
}
Sl::Gradient { x, y } => {
write!(self.out, ", ")?;
self.write_expr(module, x, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, y, func_ctx)?;
}
}
if let Some(offset) = offset {
write!(self.out, ", ")?;
self.write_const_expression(module, offset, func_ctx.expressions)?;
}
write!(self.out, ")")?;
}
Expression::ImageSample {
image,
sampler,
gather: Some(component),
coordinate,
array_index,
offset,
level: _,
depth_ref,
clamp_to_edge: _,
} => {
let suffix_cmp = match depth_ref {
Some(_) => "Compare",
None => "",
};
write!(self.out, "textureGather{suffix_cmp}(")?;
match *func_ctx.resolve_type(image, &module.types) {
TypeInner::Image {
class: crate::ImageClass::Depth { multi: _ },
..
} => {}
_ => {
write!(self.out, "{}, ", component as u8)?;
}
}
self.write_expr(module, image, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, sampler, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, coordinate, func_ctx)?;
if let Some(array_index) = array_index {
write!(self.out, ", ")?;
self.write_expr(module, array_index, func_ctx)?;
}
if let Some(depth_ref) = depth_ref {
write!(self.out, ", ")?;
self.write_expr(module, depth_ref, func_ctx)?;
}
if let Some(offset) = offset {
write!(self.out, ", ")?;
self.write_const_expression(module, offset, func_ctx.expressions)?;
}
write!(self.out, ")")?;
}
Expression::ImageQuery { image, query } => {
use crate::ImageQuery as Iq;
let texture_function = match query {
Iq::Size { .. } => "textureDimensions",
Iq::NumLevels => "textureNumLevels",
Iq::NumLayers => "textureNumLayers",
Iq::NumSamples => "textureNumSamples",
};
write!(self.out, "{texture_function}(")?;
self.write_expr(module, image, func_ctx)?;
if let Iq::Size { level: Some(level) } = query {
write!(self.out, ", ")?;
self.write_expr(module, level, func_ctx)?;
};
write!(self.out, ")")?;
}
Expression::ImageLoad {
image,
coordinate,
array_index,
sample,
level,
} => {
write!(self.out, "textureLoad(")?;
self.write_expr(module, image, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, coordinate, func_ctx)?;
if let Some(array_index) = array_index {
write!(self.out, ", ")?;
self.write_expr(module, array_index, func_ctx)?;
}
if let Some(index) = sample.or(level) {
write!(self.out, ", ")?;
self.write_expr(module, index, func_ctx)?;
}
write!(self.out, ")")?;
}
Expression::GlobalVariable(handle) => {
let name = &self.names[&NameKey::GlobalVariable(handle)];
write!(self.out, "{name}")?;
}
Expression::As {
expr,
kind,
convert,
} => {
let inner = func_ctx.resolve_type(expr, &module.types);
match *inner {
TypeInner::Matrix {
columns,
rows,
scalar,
} => {
let scalar = crate::Scalar {
kind,
width: convert.unwrap_or(scalar.width),
};
let scalar_kind_str = scalar.to_wgsl_if_implemented()?;
write!(
self.out,
"mat{}x{}<{}>",
common::vector_size_str(columns),
common::vector_size_str(rows),
scalar_kind_str
)?;
}
TypeInner::Vector {
size,
scalar: crate::Scalar { width, .. },
} => {
let scalar = crate::Scalar {
kind,
width: convert.unwrap_or(width),
};
let vector_size_str = common::vector_size_str(size);
let scalar_kind_str = scalar.to_wgsl_if_implemented()?;
if convert.is_some() {
write!(self.out, "vec{vector_size_str}<{scalar_kind_str}>")?;
} else {
write!(self.out, "bitcast<vec{vector_size_str}<{scalar_kind_str}>>")?;
}
}
TypeInner::Scalar(crate::Scalar { width, .. }) => {
let scalar = crate::Scalar {
kind,
width: convert.unwrap_or(width),
};
let scalar_kind_str = scalar.to_wgsl_if_implemented()?;
if convert.is_some() {
write!(self.out, "{scalar_kind_str}")?
} else {
write!(self.out, "bitcast<{scalar_kind_str}>")?
}
}
_ => {
return Err(Error::Unimplemented(format!(
"write_expr expression::as {inner:?}"
)));
}
};
write!(self.out, "(")?;
self.write_expr(module, expr, func_ctx)?;
write!(self.out, ")")?;
}
Expression::Load { pointer } => {
let is_atomic_pointer = func_ctx
.resolve_type(pointer, &module.types)
.is_atomic_pointer(&module.types);
if is_atomic_pointer {
write!(self.out, "atomicLoad(")?;
self.write_expr(module, pointer, func_ctx)?;
write!(self.out, ")")?;
} else {
self.write_expr_with_indirection(
module,
pointer,
func_ctx,
Indirection::Reference,
)?;
}
}
Expression::LocalVariable(handle) => {
write!(self.out, "{}", self.names[&func_ctx.name_key(handle)])?
}
Expression::ArrayLength(expr) => {
write!(self.out, "arrayLength(")?;
self.write_expr(module, expr, func_ctx)?;
write!(self.out, ")")?;
}
Expression::Math {
fun,
arg,
arg1,
arg2,
arg3,
} => {
use crate::MathFunction as Mf;
enum Function {
Regular(&'static str),
InversePolyfill(InversePolyfill),
}
let function = match fun.try_to_wgsl() {
Some(name) => Function::Regular(name),
None => match fun {
Mf::Inverse => {
let ty = func_ctx.resolve_type(arg, &module.types);
let Some(overload) = InversePolyfill::find_overload(ty) else {
return Err(Error::unsupported("math function", fun));
};
Function::InversePolyfill(overload)
}
_ => return Err(Error::unsupported("math function", fun)),
},
};
match function {
Function::Regular(fun_name) => {
write!(self.out, "{fun_name}(")?;
self.write_expr(module, arg, func_ctx)?;
for arg in IntoIterator::into_iter([arg1, arg2, arg3]).flatten() {
write!(self.out, ", ")?;
self.write_expr(module, arg, func_ctx)?;
}
write!(self.out, ")")?
}
Function::InversePolyfill(inverse) => {
write!(self.out, "{}(", inverse.fun_name)?;
self.write_expr(module, arg, func_ctx)?;
write!(self.out, ")")?;
self.required_polyfills.insert(inverse);
}
}
}
Expression::Swizzle {
size,
vector,
pattern,
} => {
self.write_expr(module, vector, func_ctx)?;
write!(self.out, ".")?;
for &sc in pattern[..size as usize].iter() {
self.out.write_char(back::COMPONENTS[sc as usize])?;
}
}
Expression::Unary { op, expr } => {
let unary = match op {
crate::UnaryOperator::Negate => "-",
crate::UnaryOperator::LogicalNot => "!",
crate::UnaryOperator::BitwiseNot => "~",
};
write!(self.out, "{unary}(")?;
self.write_expr(module, expr, func_ctx)?;
write!(self.out, ")")?
}
Expression::Select {
condition,
accept,
reject,
} => {
write!(self.out, "select(")?;
self.write_expr(module, reject, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, accept, func_ctx)?;
write!(self.out, ", ")?;
self.write_expr(module, condition, func_ctx)?;
write!(self.out, ")")?
}
Expression::Derivative { axis, ctrl, expr } => {
use crate::{DerivativeAxis as Axis, DerivativeControl as Ctrl};
let op = match (axis, ctrl) {
(Axis::X, Ctrl::Coarse) => "dpdxCoarse",
(Axis::X, Ctrl::Fine) => "dpdxFine",
(Axis::X, Ctrl::None) => "dpdx",
(Axis::Y, Ctrl::Coarse) => "dpdyCoarse",
(Axis::Y, Ctrl::Fine) => "dpdyFine",
(Axis::Y, Ctrl::None) => "dpdy",
(Axis::Width, Ctrl::Coarse) => "fwidthCoarse",
(Axis::Width, Ctrl::Fine) => "fwidthFine",
(Axis::Width, Ctrl::None) => "fwidth",
};
write!(self.out, "{op}(")?;
self.write_expr(module, expr, func_ctx)?;
write!(self.out, ")")?
}
Expression::Relational { fun, argument } => {
use crate::RelationalFunction as Rf;
let fun_name = match fun {
Rf::All => "all",
Rf::Any => "any",
_ => return Err(Error::UnsupportedRelationalFunction(fun)),
};
write!(self.out, "{fun_name}(")?;
self.write_expr(module, argument, func_ctx)?;
write!(self.out, ")")?
}
// Not supported yet
Expression::RayQueryGetIntersection { .. }
| Expression::RayQueryVertexPositions { .. } => unreachable!(),
// Nothing to do here, since call expression already cached
Expression::CallResult(_)
| Expression::AtomicResult { .. }
| Expression::RayQueryProceedResult
| Expression::SubgroupBallotResult
| Expression::SubgroupOperationResult { .. }
| Expression::WorkGroupUniformLoadResult { .. } => {}
}
Ok(())
}
/// Helper method used to write global variables
/// # Notes
/// Always adds a newline
fn write_global(
&mut self,
module: &Module,
global: &crate::GlobalVariable,
handle: Handle<crate::GlobalVariable>,
) -> BackendResult {
// Write group and binding attributes if present
if let Some(ref binding) = global.binding {
self.write_attributes(&[
Attribute::Group(binding.group),
Attribute::Binding(binding.binding),
])?;
writeln!(self.out)?;
}
// First write global name and address space if supported
write!(self.out, "var")?;
let (address, maybe_access) = address_space_str(global.space);
if let Some(space) = address {
write!(self.out, "<{space}")?;
if let Some(access) = maybe_access {
write!(self.out, ", {access}")?;
}
write!(self.out, ">")?;
}
write!(
self.out,
" {}: ",
&self.names[&NameKey::GlobalVariable(handle)]
)?;
// Write global type
self.write_type(module, global.ty)?;
// Write initializer
if let Some(init) = global.init {
write!(self.out, " = ")?;
self.write_const_expression(module, init, &module.global_expressions)?;
}
// End with semicolon
writeln!(self.out, ";")?;
Ok(())
}
/// Helper method used to write global constants
///
/// # Notes
/// Ends in a newline
fn write_global_constant(
&mut self,
module: &Module,
handle: Handle<crate::Constant>,
) -> BackendResult {
let name = &self.names[&NameKey::Constant(handle)];
// First write only constant name
write!(self.out, "const {name}: ")?;
self.write_type(module, module.constants[handle].ty)?;
write!(self.out, " = ")?;
let init = module.constants[handle].init;
self.write_const_expression(module, init, &module.global_expressions)?;
writeln!(self.out, ";")?;
Ok(())
}
#[allow(clippy::missing_const_for_fn)]
pub fn finish(self) -> W {
self.out
}
}
struct WriterTypeContext<'m> {
module: &'m Module,
names: &'m crate::FastHashMap<NameKey, String>,
}
impl TypeContext for WriterTypeContext<'_> {
fn lookup_type(&self, handle: Handle<crate::Type>) -> &crate::Type {
&self.module.types[handle]
}
fn type_name(&self, handle: Handle<crate::Type>) -> &str {
self.names[&NameKey::Type(handle)].as_str()
}
fn write_unnamed_struct<W: Write>(&self, _: &TypeInner, _: &mut W) -> core::fmt::Result {
unreachable!("the WGSL back end should always provide type handles");
}
fn write_override<W: Write>(&self, _: Handle<crate::Override>, _: &mut W) -> core::fmt::Result {
unreachable!("overrides should be validated out");
}
fn write_non_wgsl_inner<W: Write>(&self, _: &TypeInner, _: &mut W) -> core::fmt::Result {
unreachable!("backends should only be passed validated modules");
}
fn write_non_wgsl_scalar<W: Write>(&self, _: crate::Scalar, _: &mut W) -> core::fmt::Result {
unreachable!("backends should only be passed validated modules");
}
}
fn map_binding_to_attribute(binding: &crate::Binding) -> Vec<Attribute> {
match *binding {
crate::Binding::BuiltIn(built_in) => {
if let crate::BuiltIn::Position { invariant: true } = built_in {
vec![Attribute::BuiltIn(built_in), Attribute::Invariant]
} else {
vec![Attribute::BuiltIn(built_in)]
}
}
crate::Binding::Location {
location,
interpolation,
sampling,
blend_src: None,
} => vec![
Attribute::Location(location),
Attribute::Interpolate(interpolation, sampling),
],
crate::Binding::Location {
location,
interpolation,
sampling,
blend_src: Some(blend_src),
} => vec![
Attribute::Location(location),
Attribute::BlendSrc(blend_src),
Attribute::Interpolate(interpolation, sampling),
],
}
}