/// 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)] pubstruct WriterFlags: u32 { /// Always annotate the type information instead of inferring. const EXPLICIT_TYPES = 0x1;
}
}
fn reset(&mutself, module: &Module) { self.names.clear(); self.namer.reset(
module,
&crate::keywords::wgsl::RESERVED_SET,
&crate::keywords::wgsl::BUILTIN_IDENTIFIER_SET, // an identifier must not start with two underscore
proc::CaseInsensitiveKeywordSet::empty(),
&["__", "_naga"],
&mutself.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. /// /// [predeclared]: https://www.w3.org/TR/WGSL/#predeclared /// [`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
}
// 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::Mesh => { let mesh_output_name = module.global_variables
[ep.mesh_info.as_ref().unwrap().output_variable]
.name
.clone()
.unwrap(); letmut mesh_attrs = vec![
Attribute::MeshStage(mesh_output_name),
Attribute::WorkGroupSize(ep.workgroup_size),
]; iflet Some(task_payload) = ep.task_payload { let payload_name =
module.global_variables[task_payload].name.clone().unwrap();
mesh_attrs.push(Attribute::TaskPayload(payload_name));
}
mesh_attrs
}
ShaderStage::Task => { let payload_name = module.global_variables[ep.task_payload.unwrap()]
.name
.clone()
.unwrap();
vec![
Attribute::Stage(ShaderStage::Task),
Attribute::TaskPayload(payload_name),
Attribute::WorkGroupSize(ep.workgroup_size),
]
}
ShaderStage::RayGeneration => vec![Attribute::Stage(ShaderStage::RayGeneration)],
ShaderStage::AnyHit | ShaderStage::ClosestHit | ShaderStage::Miss => { let payload_name = module.global_variables[ep.incoming_ray_payload.unwrap()]
.name
.clone()
.unwrap();
vec![
Attribute::Stage(ep.stage),
Attribute::IncomingRayPayload(payload_name),
]
}
}; 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)?;
}
for ep in &module.entry_points { iflet Some(res) = ep.function.result.as_ref().and_then(|a| a.binding.as_ref()) {
check_binding(res, &mut needed);
} for arg in ep
.function
.arguments
.iter()
.filter_map(|a| a.binding.as_ref())
{
check_binding(arg, &mut needed);
}
}
// Write required declarations letmut any_written = false; if needed.f16 {
writeln!(self.out, "enable f16;")?;
any_written = true;
} if needed.int16 {
writeln!(self.out, "enable wgpu_int16;")?;
any_written = true;
} if needed.dual_source_blending {
writeln!(self.out, "enable dual_source_blending;")?;
any_written = true;
} if needed.clip_distances {
writeln!(self.out, "enable clip_distances;")?;
any_written = true;
} if module.uses_mesh_shaders() {
writeln!(self.out, "enable wgpu_mesh_shader;")?;
any_written = true;
} if needed.binding_array {
writeln!(self.out, "enable wgpu_binding_array;")?;
any_written = true;
} if needed.draw_index {
writeln!(self.out, "enable draw_index;")?;
any_written = true;
} if needed.primitive_index {
writeln!(self.out, "enable primitive_index;")?;
any_written = true;
} if needed.cooperative_matrix {
writeln!(self.out, "enable wgpu_cooperative_matrix;")?;
any_written = true;
} if needed.ray_tracing_pipeline {
writeln!(self.out, "enable wgpu_ray_tracing_pipeline;")?;
any_written = true;
} if needed.per_vertex {
writeln!(self.out, "enable wgpu_per_vertex;")?;
any_written = true;
} if any_written { // Empty line for readability
writeln!(self.out)?;
}
Ok(())
}
/// Helper method used to write /// [functions](https://gpuweb.github.io/gpuweb/wgsl/#functions) /// /// # Notes /// Ends in a newline fn write_function(
&mutself,
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 iflet 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 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 iflet 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(&mutself, 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 => "task", //Handled by another variant in the Attribute enum, so this code should never be hit.
ShaderStage::Mesh => unreachable!(),
ShaderStage::RayGeneration => "ray_generation",
ShaderStage::AnyHit => "any_hit",
ShaderStage::ClosestHit => "closest_hit",
ShaderStage::Miss => "miss",
};
/// 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(
&mutself,
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)?; iflet 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(&mutself, 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, &mutself.out)?;
Ok(())
}
fn write_type_resolution(
&mutself,
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, &mutself.out)?;
match *stmt {
Statement::Emit(ref range) => { for handle in range.clone() { let info = &func_ctx.info[handle]; let expr_name = iflet 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
}
};
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")?; iflet 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);
let l2 = level.next(); letmut 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(),
));
}
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 iflet 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, ";")?;
}
/// 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 { usecrate::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. ifself.named_expressions.contains_key(&expr) { return Indirection::Ordinary;
}
/// Write the ordinary WGSL form of `expr`. /// /// See `write_expr_with_indirection` for details. fn write_expr(
&mutself,
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(
&mutself,
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);
log::trace!( "expression {:?}={:?} is {:?}, expected {:?}",
expr,
func_ctx.expressions[expr],
plain,
requested,
); 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)?,
}
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::U16(value) => write!(self.out, "u16({value})")?, crate::Literal::I16(value) => write!(self.out, "i16({value})")?, 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, ")")?;
}
Expression::Override(handle) => {
write!(self.out, "{}", self.names[&NameKey::Override(handle)])?;
}
_ => 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(
&mutself,
module: &Module,
expr: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
indirection: Indirection,
) -> BackendResult { usecrate::Expression;
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();
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));
};
/// Helper method used to write global variables /// # Notes /// Always adds a newline fn write_global(
&mutself,
module: &Module,
global: &crate::GlobalVariable,
handle: Handle<crate::GlobalVariable>,
) -> BackendResult { // Write group and binding attributes if present iflet Some(ref binding) = global.binding { self.write_attributes(&[
Attribute::Group(binding.group),
Attribute::Binding(binding.binding),
])?;
writeln!(self.out)?;
}
if global
.memory_decorations
.contains(crate::MemoryDecorations::COHERENT)
{
write!(self.out, "@coherent ")?;
} if global
.memory_decorations
.contains(crate::MemoryDecorations::VOLATILE)
{
write!(self.out, "@volatile ")?;
}
// First write global name and address space if supported
write!(self.out, "var")?; let (address, maybe_access) = address_space_str(global.space); iflet Some(space) = address {
write!(self.out, "<{space}")?; iflet 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)?;
/// Helper method used to write global constants /// /// # Notes /// Ends in a newline fn write_global_constant(
&mutself,
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(())
}
/// Helper method used to write overrides /// /// # Notes /// Ends in a newline fn write_override(
&mutself,
module: &Module,
handle: Handle<crate::Override>,
) -> BackendResult { let override_ = &module.overrides[handle]; let name = &self.names[&NameKey::Override(handle)];
fn write_unnamed_struct<W: Write>(&self, _: &TypeInner, _: &mut W) -> core::fmt::Result {
unreachable!("the WGSL back end should always provide type handles");
}
Die Informationen auf dieser Webseite wurden
nach bestem Wissen sorgfältig zusammengestellt. Es wird jedoch weder Vollständigkeit, noch Richtigkeit,
noch Qualität der bereit gestellten Informationen zugesichert.
Bemerkung:
Die farbliche Syntaxdarstellung und die Messung sind noch experimentell.
