use alloc::{
format,
string::{String, ToString},
vec,
vec::Vec,
}; use core::{
cmp::Ordering,
fmt::{Display, Error as FmtError, Formatter, Write},
iter,
}; use num_traits::real::Real as _;
// This is a hack: we need to pass a pointer to an atomic, // but generally the backend isn't putting "&" in front of every pointer. // Some more general handling of pointers is needed to be implemented here. const ATOMIC_REFERENCE: &str = "&";
pub(crate) const ATOMIC_COMP_EXCH_FUNCTION: &str = "naga_atomic_compare_exchange_weak_explicit"; pub(crate) const MODF_FUNCTION: &str = "naga_modf"; pub(crate) const FREXP_FUNCTION: &str = "naga_frexp"; pub(crate) const ABS_FUNCTION: &str = "naga_abs"; pub(crate) const DIV_FUNCTION: &str = "naga_div"; pub(crate) const DOT_FUNCTION_PREFIX: &str = "naga_dot"; pub(crate) const MOD_FUNCTION: &str = "naga_mod"; pub(crate) const NEG_FUNCTION: &str = "naga_neg"; pub(crate) const F2I32_FUNCTION: &str = "naga_f2i32"; pub(crate) const F2U32_FUNCTION: &str = "naga_f2u32"; pub(crate) const F2I64_FUNCTION: &str = "naga_f2i64"; pub(crate) const F2U64_FUNCTION: &str = "naga_f2u64"; pub(crate) const IMAGE_LOAD_EXTERNAL_FUNCTION: &str = "nagaTextureLoadExternal"; pub(crate) const IMAGE_SIZE_EXTERNAL_FUNCTION: &str = "nagaTextureDimensionsExternal"; pub(crate) const IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION: &str = "nagaTextureSampleBaseClampToEdge"; /// For some reason, Metal does not let you have `metal::texture<..>*` as a buffer argument. /// However, if you put that texture inside a struct, everything is totally fine. This /// baffles me to no end. /// /// As such, we wrap all argument buffers in a struct that has a single generic `<T>` field. /// This allows `NagaArgumentBufferWrapper<metal::texture<..>>*` to work. The astute among /// you have noticed that this should be exactly the same to the compiler, and you're correct. pub(crate) const ARGUMENT_BUFFER_WRAPPER_STRUCT: &str = "NagaArgumentBufferWrapper"; /// Name of the struct that is declared to wrap the 3 textures and parameters /// buffer that [`crate::ImageClass::External`] variables are lowered to, /// allowing them to be conveniently passed to user-defined or wrapper /// functions. The struct is declared in [`Writer::write_type_defs`]. pub(crate) const EXTERNAL_TEXTURE_WRAPPER_STRUCT: &str = "NagaExternalTextureWrapper"; pub(crate) const COOPERATIVE_LOAD_FUNCTION: &str = "NagaCooperativeLoad"; pub(crate) const COOPERATIVE_MULTIPLY_ADD_FUNCTION: &str = "NagaCooperativeMultiplyAdd";
/// Write the Metal name for a Naga numeric type: scalar, vector, or matrix. /// /// The `sizes` slice determines whether this function writes a /// scalar, vector, or matrix type: /// /// - An empty slice produces a scalar type. /// - A one-element slice produces a vector type. /// - A two element slice `[ROWS COLUMNS]` produces a matrix of the given size. fn put_numeric_type(
out: &mutimpl Write,
scalar: crate::Scalar,
sizes: &[crate::VectorSize],
) -> Result<(), FmtError> { match (scalar, sizes) {
(scalar, &[]) => {
write!(out, "{}", scalar.to_msl_name())
}
(scalar, &[rows]) => {
write!(
out, "{}::{}{}",
NAMESPACE,
scalar.to_msl_name(),
common::vector_size_str(rows)
)
}
(scalar, &[rows, columns]) => {
write!(
out, "{}::{}{}x{}",
NAMESPACE,
scalar.to_msl_name(),
common::vector_size_str(columns),
common::vector_size_str(rows)
)
}
(_, _) => Ok(()), // not meaningful
}
}
/// Prefix for cached clamped level-of-detail values for `ImageLoad` expressions. const CLAMPED_LOD_LOAD_PREFIX: &str = "clamped_lod_e";
/// Prefix for reinterpreted expressions using `as_type<T>(...)`. const REINTERPRET_PREFIX: &str = "reinterpreted_";
/// Wrapper for identifier names for clamped level-of-detail values /// /// Values of this type implement [`core::fmt::Display`], formatting as /// the name of the variable used to hold the cached clamped /// level-of-detail value for an `ImageLoad` expression. struct ClampedLod(Handle<crate::Expression>);
/// Wrapper for generating `struct _mslBufferSizes` member names for /// runtime-sized array lengths. /// /// On Metal, `wgpu_hal` passes the element counts for all runtime-sized arrays /// as an argument to the entry point. This argument's type in the MSL is /// `struct _mslBufferSizes`, a Naga-synthesized struct with a `uint` member for /// each global variable containing a runtime-sized array. /// /// If `global` is a [`Handle`] for a [`GlobalVariable`] that contains a /// runtime-sized array, then the value `ArraySize(global)` implements /// [`core::fmt::Display`], formatting as the name of the struct member carrying /// the number of elements in that runtime-sized array. /// /// [`GlobalVariable`]: crate::GlobalVariable struct ArraySizeMember(Handle<crate::GlobalVariable>);
/// Wrapper for reinterpreted variables using `as_type<target_type>(orig)`. /// /// Implements [`core::fmt::Display`], formatting as a name derived from /// `target_type` and the variable name of `orig`. #[derive(Clone, Copy)] struct Reinterpreted<'a> {
target_type: &'a str,
orig: Handle<crate::Expression>,
}
impl TypedGlobalVariable<'_> { fn to_parts(&self) -> Result<TypedGlobalVariableParts, Error> { let var = &self.module.global_variables[self.handle]; let name = &self.names[&NameKey::GlobalVariable(self.handle)];
pubstruct Writer<W> { pub(super) out: W, pub(super) names: FastHashMap<NameKey, String>, pub(super) named_expressions: crate::NamedExpressions, /// Set of expressions that need to be baked to avoid unnecessary repetition in output
need_bake_expressions: back::NeedBakeExpressions, pub(super) namer: proc::Namer, pub(super) wrapped_functions: FastHashSet<WrappedFunction>, #[cfg(test)]
put_expression_stack_pointers: FastHashSet<*const ()>, #[cfg(test)]
put_block_stack_pointers: FastHashSet<*const ()>, /// Set of (struct type, struct field index) denoting which fields require /// padding inserted **before** them (i.e. between fields at index - 1 and index)
struct_member_pads: FastHashSet<(Handle<crate::Type>, u32)>,
needs_object_memory_barriers: bool,
}
implcrate::AddressSpace { /// Returns true if global variables in this address space are /// passed in function arguments. These arguments need to be /// passed through any functions called from the entry point. constfn needs_pass_through(&self) -> bool { match *self { Self::Uniform
| Self::Storage { .. }
| Self::Private
| Self::WorkGroup
| Self::Immediate
| Self::Handle
| Self::TaskPayload => true, Self::Function => false, Self::RayPayload | Self::IncomingRayPayload => unreachable!(),
}
}
/// Returns true if the address space may need a "const" qualifier. constfn needs_access_qualifier(&self) -> bool { match *self { //Note: we are ignoring the storage access here, and instead // rely on the actual use of a global by functions. This means we // may end up with "const" even if the binding is read-write, // and that should be OK. Self::Storage { .. } => true, Self::TaskPayload => true, Self::RayPayload | Self::IncomingRayPayload => unimplemented!(), // These should always be read-write. Self::Private | Self::WorkGroup => false, // These translate to `constant` address space, no need for qualifiers. Self::Uniform | Self::Immediate => false, // Not applicable. Self::Handle | Self::Function => false,
}
}
constfn to_msl_name(self) -> Option<&'static str> { matchself { Self::Handle => None, Self::Uniform | Self::Immediate => Some("constant"), Self::Storage { .. } => Some("device"), // note for `RayPayload`, this probably needs to be emulated as a // private variable, as metal has essentially an inout input // for where it is passed. Self::Private | Self::Function | Self::RayPayload => Some("thread"), Self::WorkGroup => Some("threadgroup"), Self::TaskPayload => Some("object_data"), Self::IncomingRayPayload => Some("ray_data"),
}
}
}
implcrate::Type { // Returns `true` if we need to emit an alias for this type. constfn needs_alias(&self) -> bool { usecrate::TypeInner as Ti;
matchself.inner { // value types are concise enough, we only alias them if they are named
Ti::Scalar(_)
| Ti::Vector { .. }
| Ti::Matrix { .. }
| Ti::CooperativeMatrix { .. }
| Ti::Atomic(_)
| Ti::Pointer { .. }
| Ti::ValuePointer { .. } => self.name.is_some(), // composite types are better to be aliased, regardless of the name
Ti::Struct { .. } | Ti::Array { .. } => true, // handle types may be different, depending on the global var access, so we always inline them
Ti::Image { .. }
| Ti::Sampler { .. }
| Ti::AccelerationStructure { .. }
| Ti::RayQuery { .. }
| Ti::BindingArray { .. } => false,
}
}
}
/// Return the name key for a local variable used by ReadZeroSkipWrite bounds-check /// policy when it needs to produce a pointer-typed result for an OOB access. These /// are unique per accessed type, so the second argument is a type handle. See docs /// for [`crate::back::msl`]. fn oob_local_for_type(origin: FunctionOrigin, ty: Handle<crate::Type>) -> NameKey { match origin {
FunctionOrigin::Handle(handle) => NameKey::FunctionOobLocal(handle, ty),
FunctionOrigin::EntryPoint(idx) => NameKey::EntryPointOobLocal(idx, ty),
}
}
}
impl NameKeyExt for NameKey {}
/// A level of detail argument. /// /// When [`BoundsCheckPolicy::Restrict`] applies to an [`ImageLoad`] access, we /// save the clamped level of detail in a temporary variable whose name is based /// on the handle of the `ImageLoad` expression. But for other policies, we just /// use the expression directly. /// /// [`BoundsCheckPolicy::Restrict`]: index::BoundsCheckPolicy::Restrict /// [`ImageLoad`]: crate::Expression::ImageLoad #[derive(Clone, Copy)] enum LevelOfDetail {
Direct(Handle<crate::Expression>),
Restricted(Handle<crate::Expression>),
}
/// Values needed to select a particular texel for [`ImageLoad`] and [`ImageStore`]. /// /// When this is used in code paths unconcerned with the `Restrict` bounds check /// policy, the `LevelOfDetail` enum introduces an unneeded match, since `level` /// will always be either `None` or `Some(Direct(_))`. But this turns out not to /// be too awkward. If that changes, we can revisit. /// /// [`ImageLoad`]: crate::Expression::ImageLoad /// [`ImageStore`]: crate::Statement::ImageStore struct TexelAddress {
coordinate: Handle<crate::Expression>,
array_index: Option<Handle<crate::Expression>>,
sample: Option<Handle<crate::Expression>>,
level: Option<LevelOfDetail>,
}
/// The set of expressions used as indices in `ReadZeroSkipWrite`-policy /// accesses. These may need to be cached in temporary variables. See /// `index::find_checked_indexes` for details. pub(super) guarded_indices: HandleSet<crate::Expression>, /// See [`Writer::gen_force_bounded_loop_statements`] for details. pub(super) force_loop_bounding: bool,
}
/// Return true if calls to `image`'s `read` and `write` methods should supply a level of detail. /// /// Only mipmapped images need to specify a level of detail. Since 1D /// textures cannot have mipmaps, MSL requires that the level argument to /// texture1d queries and accesses must be a constexpr 0. It's easiest /// just to omit the level entirely for 1D textures. fn image_needs_lod(&self, image: Handle<crate::Expression>) -> bool { let image_ty = self.resolve_type(image); ifletcrate::TypeInner::Image { dim, class, .. } = *image_ty {
class.is_mipmapped() && dim != crate::ImageDimension::D1
} else { false
}
}
/// Generates statements to be inserted immediately before and at the very /// start of the body of each loop, to defeat MSL infinite loop reasoning. /// The 0th item of the returned tuple should be inserted immediately prior /// to the loop and the 1st item should be inserted at the very start of /// the loop body. /// /// # What is this trying to solve? /// /// In Metal Shading Language, an infinite loop has undefined behavior. /// (This rule is inherited from C++14.) This means that, if the MSL /// compiler determines that a given loop will never exit, it may assume /// that it is never reached. It may thus assume that any conditions /// sufficient to cause the loop to be reached must be false. Like many /// optimizing compilers, MSL uses this kind of analysis to establish limits /// on the range of values variables involved in those conditions might /// hold. /// /// For example, suppose the MSL compiler sees the code: /// /// ```ignore /// if (i >= 10) { /// while (true) { } /// } /// ``` /// /// It will recognize that the `while` loop will never terminate, conclude /// that it must be unreachable, and thus infer that, if this code is /// reached, then `i < 10` at that point. /// /// Now suppose that, at some point where `i` has the same value as above, /// the compiler sees the code: /// /// ```ignore /// if (i < 10) { /// a[i] = 1; /// } /// ``` /// /// Because the compiler is confident that `i < 10`, it will make the /// assignment to `a[i]` unconditional, rewriting this code as, simply: /// /// ```ignore /// a[i] = 1; /// ``` /// /// If that `if` condition was injected by Naga to implement a bounds check, /// the MSL compiler's optimizations could allow out-of-bounds array /// accesses to occur. /// /// Naga cannot feasibly anticipate whether the MSL compiler will determine /// that a loop is infinite, so an attacker could craft a Naga module /// containing an infinite loop protected by conditions that cause the Metal /// compiler to remove bounds checks that Naga injected elsewhere in the /// function. /// /// This rewrite could occur even if the conditional assignment appears /// *before* the `while` loop, as long as `i < 10` by the time the loop is /// reached. This would allow the attacker to save the results of /// unauthorized reads somewhere accessible before entering the infinite /// loop. But even worse, the MSL compiler has been observed to simply /// delete the infinite loop entirely, so that even code dominated by the /// loop becomes reachable. This would make the attack even more flexible, /// since shaders that would appear to never terminate would actually exit /// nicely, after having stolen data from elsewhere in the GPU address /// space. /// /// To avoid UB, Naga must persuade the MSL compiler that no loop Naga /// generates is infinite. One approach would be to add inline assembly to /// each loop that is annotated as potentially branching out of the loop, /// but which in fact generates no instructions. Unfortunately, inline /// assembly is not handled correctly by some Metal device drivers. /// /// A previously used approach was to add the following code to the bottom /// of every loop: /// /// ```ignore /// if (volatile bool unpredictable = false; unpredictable) /// break; /// ``` /// /// Although the `if` condition will always be false in any real execution, /// the `volatile` qualifier prevents the compiler from assuming this. Thus, /// it must assume that the `break` might be reached, and hence that the /// loop is not unbounded. This prevents the range analysis impact described /// above. Unfortunately this prevented the compiler from making important, /// and safe, optimizations such as loop unrolling and was observed to /// significantly hurt performance. /// /// Our current approach declares a counter before every loop and /// increments it every iteration, breaking after 2^64 iterations: /// /// ```ignore /// uint2 loop_bound = uint2(0); /// while (true) { /// if (metal::all(loop_bound == uint2(4294967295))) { break; } /// loop_bound += uint2(loop_bound.y == 4294967295, 1); /// } /// ``` /// /// This convinces the compiler that the loop is finite and therefore may /// execute, whilst at the same time allowing optimizations such as loop /// unrolling. Furthermore the 64-bit counter is large enough it seems /// implausible that it would affect the execution of any shader. /// /// This approach is also used by Chromium WebGPU's Dawn shader compiler: /// <https://dawn.googlesource.com/dawn/+/d9e2d1f718678ebee0728b999830576c410cce0a/src/tint/lang/core/ir/transform/prevent_infinite_loops.cc> fn gen_force_bounded_loop_statements(
&mutself,
level: back::Level,
context: &StatementContext,
) -> Option<(String, String)> { if !context.expression.force_loop_bounding { return None;
}
let loop_bound_name = self.namer.call("loop_bound"); // Count down from u32::MAX rather than up from 0 to avoid hang on // certain Intel drivers. See <https://github.com/gfx-rs/wgpu/issues/7319>. let decl = format!("{level}uint2 {loop_bound_name} = uint2({}u);", u32::MAX); let level = level.next(); let break_and_inc = format!( "{level}if ({NAMESPACE}::all({loop_bound_name} == uint2(0u))) {{ break; }}
{level}{loop_bound_name} -= uint2({loop_bound_name}.y == 0u, 1u);"
);
fn put_call_parameters_impl<C, E>(
&mutself,
parameters: impl Iterator<Item = Handle<crate::Expression>>,
ctx: &C,
put_expression: E,
) -> BackendResult where
E: Fn(&mutSelf, &C, Handle<crate::Expression>) -> BackendResult,
{
write!(self.out, "(")?; for (i, handle) in parameters.enumerate() { if i != 0 {
write!(self.out, ", ")?;
}
put_expression(self, ctx, handle)?;
}
write!(self.out, ")")?;
Ok(())
}
/// Writes the local variables of the given function, as well as any extra /// out-of-bounds locals that are needed. /// /// The names of the OOB locals are also added to `self.names` at the same /// time. fn put_locals(&mutself, context: &ExpressionContext) -> BackendResult { let oob_local_types = context.oob_local_types(); for &ty in oob_local_types.iter() { let name_key = NameKey::oob_local_for_type(context.origin, ty); self.names.insert(name_key, self.namer.call("oob"));
}
//Note: MSL only has separate width/height/depth queries, // so compose the result of them. let dim = match *context.resolve_type(image) { crate::TypeInner::Image { dim, .. } => dim, ref other => unreachable!("Unexpected type {:?}", other),
}; let scalar = crate::Scalar { kind, width: 4 }; let coordinate_type = scalar.to_msl_name(); match dim { crate::ImageDimension::D1 => { // Since 1D textures never have mipmaps, MSL requires that the // `level` argument be a constexpr 0. It's simplest for us just // to pass `None` and omit the level entirely. if kind == crate::ScalarKind::Uint { // No need to construct a vector. No cast needed. self.put_image_query(image, "width", None, context)?;
} else { // There's no definition for `int` in the `metal` namespace.
write!(self.out, "int(")?; self.put_image_query(image, "width", None, context)?;
write!(self.out, ")")?;
}
} crate::ImageDimension::D2 => {
write!(self.out, "{NAMESPACE}::{coordinate_type}2(")?; self.put_image_query(image, "width", level, context)?;
write!(self.out, ", ")?; self.put_image_query(image, "height", level, context)?;
write!(self.out, ")")?;
} crate::ImageDimension::D3 => {
write!(self.out, "{NAMESPACE}::{coordinate_type}3(")?; self.put_image_query(image, "width", level, context)?;
write!(self.out, ", ")?; self.put_image_query(image, "height", level, context)?;
write!(self.out, ", ")?; self.put_image_query(image, "depth", level, context)?;
write!(self.out, ")")?;
} crate::ImageDimension::Cube => {
write!(self.out, "{NAMESPACE}::{coordinate_type}2(")?; self.put_image_query(image, "width", level, context)?;
write!(self.out, ")")?;
}
}
Ok(())
}
fn put_cast_to_uint_scalar_or_vector(
&mutself,
expr: Handle<crate::Expression>,
context: &ExpressionContext,
) -> BackendResult { // coordinates in IR are int, but Metal expects uint match *context.resolve_type(expr) { crate::TypeInner::Scalar(_) => {
put_numeric_type(&mutself.out, crate::Scalar::U32, &[])?
} crate::TypeInner::Vector { size, .. } => {
put_numeric_type(&mutself.out, crate::Scalar::U32, &[size])?
}
_ => { return Err(Error::GenericValidation( "Invalid type for image coordinate".into(),
))
}
};
// Write the array index, if present. iflet Some(array_index) = address.array_index {
write!(self.out, ", ")?; self.put_restricted_scalar_image_index(image, array_index, "get_array_size", context)?;
}
// Write the sample index, if present. iflet Some(sample) = address.sample {
write!(self.out, ", ")?; self.put_restricted_scalar_image_index(image, sample, "get_num_samples", context)?;
}
// The level of detail should be clamped and cached by // `put_cache_restricted_level`, so we don't need to clamp it here. iflet Some(level) = address.level {
write!(self.out, ", ")?; self.put_level_of_detail(level, context)?;
}
Ok(())
}
/// Write an expression that is true if the given image access is in bounds. fn put_image_access_bounds_check(
&mutself,
image: Handle<crate::Expression>,
address: &TexelAddress,
context: &ExpressionContext,
) -> BackendResult { letmut conjunction = "";
// First, check the level of detail. Only if that is in bounds can we // use it to find the appropriate bounds for the coordinates. let level = iflet Some(level) = address.level {
write!(self.out, "uint(")?; self.put_level_of_detail(level, context)?;
write!(self.out, ") < ")?; self.put_expression(image, context, true)?;
write!(self.out, ".get_num_mip_levels()")?;
conjunction = " && ";
Some(level)
} else {
None
};
match context.policies.image_load {
proc::BoundsCheckPolicy::Restrict => { // Use the cached restricted level of detail, if any. Omit the // level altogether for 1D textures. if address.level.is_some() {
address.level = if context.image_needs_lod(image) {
Some(LevelOfDetail::Restricted(load))
} else {
None
}
}
fn put_unchecked_image_load(
&mut self,
image: Handle<crate::Expression>,
address: &TexelAddress,
context: &ExpressionContext,
) -> BackendResult {
self.put_expression(image, context, false)?;
write!(self.out, ".read(")?; // coordinates in IR are int, but Metal expects uint
self.put_cast_to_uint_scalar_or_vector(address.coordinate, context)?; if let Some(expr) = address.array_index {
write!(self.out, ", ")?;
self.put_expression(expr, context, true)?;
} if let Some(sample) = address.sample {
write!(self.out, ", ")?;
self.put_expression(sample, context, true)?;
} if let Some(level) = address.level { if context.image_needs_lod(image) {
write!(self.out, ", ")?;
self.put_level_of_detail(level, context)?;
}
}
write!(self.out, ")")?;
Ok(())
}
fn put_image_atomic(
&mut self,
level: back::Level,
image: Handle<crate::Expression>,
address: &TexelAddress,
fun: crate::AtomicFunction,
value: Handle<crate::Expression>,
context: &StatementContext,
) -> BackendResult {
write!(self.out, "{level}")?;
self.put_expression(image, &context.expression, false)?;
let op = if context.expression.resolve_type(value).scalar_width() == Some(8) {
fun.to_msl_64_bit()?
} else {
fun.to_msl()
};
write!(self.out, ".atomic_{op}(")?; // coordinates in IR are int, but Metal expects uint
self.put_cast_to_uint_scalar_or_vector(address.coordinate, &context.expression)?;
write!(self.out, ", ")?;
self.put_expression(value, &context.expression, true)?;
writeln!(self.out, ");")?;
// Workaround for Apple Metal TBDR driver bug: fragment shader atomic // texture writes randomly drop unless followed by a standard texture // write. Insert a dead-code write behind an unprovable condition so // the compiler emits proper memory safety barriers. // See: https://projects.blender.org/blender/blender/commit/aa95220576706122d79c91c7f5c522e6c7416425
let value_ty = context.expression.resolve_type(value);
let zero_value = match (value_ty.scalar_kind(), value_ty.scalar_width()) {
(Some(crate::ScalarKind::Sint), _) => "int4(0)",
(_, Some(8)) => "ulong4(0uL)",
_ => "uint4(0u)",
};
let coord_ty = context.expression.resolve_type(address.coordinate);
let x = if matches!(coord_ty, crate::TypeInner::Scalar(_)) { ""
} else { ".x"
};
write!(self.out, "{level}if (")?;
self.put_expression(address.coordinate, &context.expression, true)?;
write!(self.out, "{x} == -99999) {{ ")?;
self.put_expression(image, &context.expression, false)?;
write!(self.out, ".write({zero_value}, ")?;
self.put_cast_to_uint_scalar_or_vector(address.coordinate, &context.expression)?; if let Some(array_index) = address.array_index {
write!(self.out, ", ")?;
self.put_expression(array_index, &context.expression, true)?;
}
writeln!(self.out, "); }}")?;
Ok(())
}
fn put_image_store(
&mut self,
level: back::Level,
image: Handle<crate::Expression>,
address: &TexelAddress,
value: Handle<crate::Expression>,
context: &StatementContext,
) -> BackendResult {
write!(self.out, "{level}")?;
self.put_expression(image, &context.expression, false)?;
write!(self.out, ".write(")?;
self.put_expression(value, &context.expression, true)?;
write!(self.out, ", ")?; // coordinates in IR are int, but Metal expects uint
self.put_cast_to_uint_scalar_or_vector(address.coordinate, &context.expression)?; if let Some(expr) = address.array_index {
write!(self.out, ", ")?;
self.put_expression(expr, &context.expression, true)?;
}
writeln!(self.out, ");")?;
Ok(())
}
/// Write the maximum valid index of the dynamically sized array at the end of `handle`. /// /// The 'maximum valid index' is simply one less than the array's length. /// /// This emits an expression of the form `a / b`, so the caller must /// parenthesize its output if it will be applying operators of higher /// precedence. /// /// `handle` must be the handle of a global variable whose final member is a /// dynamically sized array.
fn put_dynamic_array_max_index(
&mut self,
handle: Handle<crate::GlobalVariable>,
context: &ExpressionContext,
) -> BackendResult {
let global = &context.module.global_variables[handle];
let (offset, array_ty) = match context.module.types[global.ty].inner {
crate::TypeInner::Struct { ref members, .. } => match members.last() {
Some(&crate::StructMember { offset, ty, .. }) => (offset, ty),
None => return Err(Error::GenericValidation("Struct has no members".into())),
},
crate::TypeInner::Array {
size: crate::ArraySize::Dynamic,
..
} => (0, global.ty),
ref ty => { return Err(Error::GenericValidation(format!( "Expected type with dynamic array, got {ty:?}"
)))
}
};
// When the stride length is larger than the size, the final element's stride of // bytes would have padding following the value. But the buffer size in // `buffer_sizes.sizeN` may not include this padding - it only needs to be large // enough to hold the actual values' bytes. // // So subtract off the size to get a byte size that falls at the start or within // the final element. Then divide by the stride size, to get one less than the // length, and then add one. This works even if the buffer size does include the // stride padding, since division rounds towards zero (MSL 2.4 §6.1). It will fail // if there are zero elements in the array, but the WebGPU `validating shader binding` // rules, together with draw-time validation when `minBindingSize` is zero, // prevent that.
write!(
self.out, "(_buffer_sizes.{member} - {offset} - {size}) / {stride}",
member = ArraySizeMember(handle),
offset = offset,
size = size,
stride = stride,
)?;
Ok(())
}
/// Emit code for the arithmetic expression of the dot product. /// /// The argument `extractor` is a function that accepts a `Writer`, a vector, and /// an index. It writes out the expression for the vector component at that index.
fn put_dot_product<T: Copy>(
&mut self,
arg: T,
arg1: T,
size: usize,
extractor: impl Fn(&mut Self, T, usize) -> BackendResult,
) -> BackendResult { // Write parentheses around the dot product expression to prevent operators // with different precedences from applying earlier.
write!(self.out, "(")?;
// Cycle through all the components of the vector for index in 0..size { // Write the addition to the previous product // This will print an extra '+' at the beginning but that is fine in msl
write!(self.out, " + ")?;
extractor(self, arg, index)?;
write!(self.out, " * ")?;
extractor(self, arg1, index)?;
}
write!(self.out, ")")?;
Ok(())
}
/// Emit code for the WGSL functions `pack4x{I, U}8[Clamp]`.
fn put_pack4x8(
&mut self,
arg: Handle<crate::Expression>,
context: &ExpressionContext<'_>,
was_signed: bool,
clamp_bounds: Option<(&str, &str)>,
) -> Result<(), Error> {
let write_arg = |this: &mut Self| -> BackendResult { if let Some((min, max)) = clamp_bounds { // Clamping with scalar bounds works (component-wise) even for packed_[u]char4.
write!(this.out, "{NAMESPACE}::clamp(")?;
this.put_expression(arg, context, true)?;
write!(this.out, ", {min}, {max})")?;
} else {
this.put_expression(arg, context, true)?;
}
Ok(())
};
if context.lang_version >= (2, 1) {
let packed_type = if was_signed { "packed_char4"
} else { "packed_uchar4"
}; // Metal uses little endian byte order, which matches what WGSL expects here.
write!(self.out, "as_type<uint>({packed_type}(")?;
write_arg(self)?;
write!(self.out, "))")?;
} else { // MSL < 2.1 doesn't support `as_type` casting between packed chars and scalars. if was_signed {
write!(self.out, "uint(")?;
}
write!(self.out, "(")?;
write_arg(self)?;
write!(self.out, "[0] & 0xFF) | ((")?;
write_arg(self)?;
write!(self.out, "[1] & 0xFF) << 8) | ((")?;
write_arg(self)?;
write!(self.out, "[2] & 0xFF) << 16) | ((")?;
write_arg(self)?;
write!(self.out, "[3] & 0xFF) << 24)")?; if was_signed {
write!(self.out, ")")?;
}
}
Ok(())
}
/// Emit code for the isign expression. ///
fn put_isign(
&mut self,
arg: Handle<crate::Expression>,
context: &ExpressionContext,
) -> BackendResult {
write!(self.out, "{NAMESPACE}::select({NAMESPACE}::select(")?;
let scalar = context
.resolve_type(arg)
.scalar()
.expect("put_isign should only be called for args which have an integer scalar type")
.to_msl_name();
match context.resolve_type(arg) {
&crate::TypeInner::Vector { size, .. } => {
let size = common::vector_size_str(size);
write!(self.out, "{scalar}{size}(-1), {scalar}{size}(1)")?;
}
_ => {
write!(self.out, "{scalar}(-1), {scalar}(1)")?;
}
}
write!(self.out, ", (")?;
self.put_expression(arg, context, true)?;
write!(self.out, " > 0)), {scalar}(0), (")?;
self.put_expression(arg, context, true)?;
write!(self.out, " == 0))")?;
Ok(())
}
fn put_literal(&mut self, literal: crate::Literal) -> BackendResult {
match literal {
crate::Literal::F64(_) => { return Err(Error::CapabilityNotSupported(valid::Capabilities::FLOAT64))
}
crate::Literal::F16(value) => { if value.is_infinite() {
let sign = if value.is_sign_negative() { "-" } else { "" };
write!(self.out, "{sign}INFINITY")?;
} elseif value.is_nan() {
write!(self.out, "NAN")?;
} else {
let suffix = if value.fract() == f16::from_f32(0.0) { ".0h"
} else { "h"
};
write!(self.out, "{value}{suffix}")?;
}
}
crate::Literal::F32(value) => { if value.is_infinite() {
let sign = if value.is_sign_negative() { "-" } else { "" };
write!(self.out, "{sign}INFINITY")?;
} elseif value.is_nan() {
write!(self.out, "NAN")?;
} else {
let suffix = if value.fract() == 0.0 { ".0" } else { "" };
write!(self.out, "{value}{suffix}")?;
}
}
crate::Literal::U16(value) => {
write!(self.out, "static_cast<ushort>({value})")?;
}
crate::Literal::I16(value) => {
write!(self.out, "static_cast<short>({value})")?;
}
crate::Literal::U32(value) => {
write!(self.out, "{value}u")?;
}
crate::Literal::I32(value) => { // `-2147483648` is parsed as unary negation of positive 2147483648. // 2147483648 is too large for int32_t meaning the expression gets // promoted to a int64_t which is not our intention. Avoid this by instead // using `-2147483647 - 1`. if value == i32::MIN {
write!(self.out, "({} - 1)", value + 1)?;
} else {
write!(self.out, "{value}")?;
}
}
crate::Literal::U64(value) => {
write!(self.out, "{value}uL")?;
}
crate::Literal::I64(value) => { // `-9223372036854775808` is parsed as unary negation of positive // 9223372036854775808. 9223372036854775808 is too large for int64_t // causing Metal to emit a `-Wconstant-conversion` warning, and change the // value to `-9223372036854775808`. Which would then be negated, possibly // causing undefined behaviour. Avoid this by instead using // `-9223372036854775808L - 1L`. if value == i64::MIN {
write!(self.out, "({}L - 1L)", value + 1)?;
} else {
write!(self.out, "{value}L")?;
}
}
crate::Literal::Bool(value) => {
write!(self.out, "{value}")?;
}
crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => { return Err(Error::GenericValidation( "Unsupported abstract literal".into(),
));
}
}
Ok(())
}
#[allow(clippy::too_many_arguments)]
fn put_possibly_const_expression<C, I, E>(
&mut self,
expr_handle: Handle<crate::Expression>,
expressions: &crate::Arena<crate::Expression>,
module: &crate::Module,
mod_info: &valid::ModuleInfo,
ctx: &C,
get_expr_ty: I,
put_expression: E,
) -> BackendResult
where
I: Fn(&C, Handle<crate::Expression>) -> &TypeResolution,
E: Fn(&mut Self, &C, Handle<crate::Expression>) -> BackendResult,
{
match expressions[expr_handle] {
crate::Expression::Literal(literal) => {
self.put_literal(literal)?;
}
crate::Expression::Constant(handle) => {
let constant = &module.constants[handle]; if constant.name.is_some() {
write!(self.out, "{}", self.names[&NameKey::Constant(handle)])?;
} else {
self.put_const_expression(
constant.init,
module,
mod_info,
&module.global_expressions,
)?;
}
}
crate::Expression::ZeroValue(ty) => {
let ty_name = TypeContext {
handle: ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: false,
};
write!(self.out, "{ty_name} {{}}")?;
}
crate::Expression::Compose { ty, ref components } => {
let ty_name = TypeContext {
handle: ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: false,
};
write!(self.out, "{ty_name}")?;
match module.types[ty].inner {
crate::TypeInner::Scalar(_)
| crate::TypeInner::Vector { .. }
| crate::TypeInner::Matrix { .. } => {
self.put_call_parameters_impl(
components.iter().copied(),
ctx,
put_expression,
)?;
}
crate::TypeInner::Array { .. } => { // Naga Arrays are Metal arrays wrapped in structs, so // we need two levels of braces.
write!(self.out, " {{{{")?; for (index, &component) in components.iter().enumerate() { if index != 0 {
write!(self.out, ", ")?;
}
put_expression(self, ctx, component)?;
}
write!(self.out, "}}}}")?;
}
crate::TypeInner::Struct { .. } => {
write!(self.out, " {{")?; for (index, &component) in components.iter().enumerate() { if index != 0 {
write!(self.out, ", ")?;
} // insert padding initialization, if needed if self.struct_member_pads.contains(&(ty, index as u32)) {
write!(self.out, "{{}}, ")?;
}
put_expression(self, ctx, component)?;
}
write!(self.out, "}}")?;
}
_ => return Err(Error::UnsupportedCompose(ty)),
}
}
crate::Expression::Splat { size, value } => {
let scalar = match *get_expr_ty(ctx, value).inner_with(&module.types) {
crate::TypeInner::Scalar(scalar) => scalar,
ref ty => { return Err(Error::GenericValidation(format!( "Expected splat value type must be a scalar, got {ty:?}",
)))
}
};
put_numeric_type(&mut self.out, scalar, &[size])?;
write!(self.out, "(")?;
put_expression(self, ctx, value)?;
write!(self.out, ")")?;
}
_ => { return Err(Error::Override);
}
}
Ok(())
}
/// Emit code for the expression `expr_handle`. /// /// The `is_scoped` argument is true if the surrounding operators have the /// precedence of the comma operator, or lower. So, for example: /// /// - Pass `true` for `is_scoped` when writing function arguments, an /// expression statement, an initializer expression, or anything already /// wrapped in parenthesis. /// /// - Pass `false` if it is an operand of a `?:` operator, a `[]`, or really /// almost anything else.
pub(super) fn put_expression(
&mut self,
expr_handle: Handle<crate::Expression>,
context: &ExpressionContext,
is_scoped: bool,
) -> BackendResult { // Add to the set in order to track the stack size. #[cfg(test)]
self.put_expression_stack_pointers
.insert(ptr::from_ref(&expr_handle).cast());
if let Some(name) = self.named_expressions.get(&expr_handle) {
write!(self.out, "{name}")?; return Ok(());
}
let expression = &context.function.expressions[expr_handle];
match *expression {
crate::Expression::Literal(_)
| crate::Expression::Constant(_)
| crate::Expression::ZeroValue(_)
| crate::Expression::Compose { .. }
| crate::Expression::Splat { .. } => {
self.put_possibly_const_expression(
expr_handle,
&context.function.expressions,
context.module,
context.mod_info,
context,
|context, expr: Handle<crate::Expression>| &context.info[expr].ty,
|writer, context, expr| writer.put_expression(expr, context, true),
)?;
}
crate::Expression::Override(_) => return Err(Error::Override),
crate::Expression::Access { base, .. }
| crate::Expression::AccessIndex { base, .. } => { // This is an acceptable place to generate a `ReadZeroSkipWrite` check. // Since `put_bounds_checks` and `put_access_chain` handle an entire // access chain at a time, recursing back through `put_expression` only // for index expressions and the base object, we will never see intermediate // `Access` or `AccessIndex` expressions here.
let policy = context.choose_bounds_check_policy(base); if policy == index::BoundsCheckPolicy::ReadZeroSkipWrite
&& self.put_bounds_checks(
expr_handle,
context,
back::Level(0), if is_scoped { "" } else { "(" },
)?
{
write!(self.out, " ? ")?;
self.put_access_chain(expr_handle, policy, context)?;
write!(self.out, " : ")?;
if context.resolve_type(base).pointer_space().is_some() { // We can't just use `DefaultConstructible` if this is a pointer. // Instead, we create a dummy local variable to serve as pointer // target if the access is out of bounds.
let result_ty = context.info[expr_handle]
.ty
.inner_with(&context.module.types)
.pointer_base_type();
let result_ty_handle = match result_ty {
Some(TypeResolution::Handle(handle)) => handle,
Some(TypeResolution::Value(_)) => { // As long as the result of a pointer access expression is // passed to a function or stored in a let binding, the // type will be in the arena. If additional uses of // pointers become valid, this assumption might no longer // hold. Note that the LHS of a load or store doesn't // take this path -- there is dedicated code in `put_load` // and `put_store`.
unreachable!( "Expected type {result_ty:?} of access through pointer type {base:?} to be in the arena",
);
}
None => {
unreachable!( "Expected access through pointer type {base:?} to return a pointer, but got {result_ty:?}",
)
}
};
let name_key =
NameKey::oob_local_for_type(context.origin, result_ty_handle);
self.out.write_str(&self.names[&name_key])?;
} else {
write!(self.out, "DefaultConstructible()")?;
}
if !is_scoped {
write!(self.out, ")")?;
}
} else {
self.put_access_chain(expr_handle, policy, context)?;
}
}
crate::Expression::Swizzle {
size,
vector,
pattern,
} => {
self.put_wrapped_expression_for_packed_vec3_access(
vector,
context, false,
&Self::put_expression,
)?;
write!(self.out, ".")?; for &sc in pattern[..size as usize].iter() {
write!(self.out, "{}", back::COMPONENTS[sc as usize])?;
}
}
crate::Expression::FunctionArgument(index) => {
let name_key = match context.origin {
FunctionOrigin::Handle(handle) => NameKey::FunctionArgument(handle, index),
FunctionOrigin::EntryPoint(ep_index) => {
NameKey::EntryPointArgument(ep_index, index)
}
};
let name = &self.names[&name_key];
write!(self.out, "{name}")?;
}
crate::Expression::GlobalVariable(handle) => {
let name = &self.names[&NameKey::GlobalVariable(handle)];
write!(self.out, "{name}")?;
}
crate::Expression::LocalVariable(handle) => {
let name_key = NameKey::local(context.origin, handle);
let name = &self.names[&name_key];
write!(self.out, "{name}")?;
}
crate::Expression::Load { pointer } => self.put_load(pointer, context, is_scoped)?,
crate::Expression::ImageSample {
coordinate,
image,
sampler,
clamp_to_edge: true,
gather: None,
array_index: None,
offset: None,
level: crate::SampleLevel::Zero,
depth_ref: None,
} => {
write!(self.out, "{IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION}(")?;
self.put_expression(image, context, true)?;
write!(self.out, ", ")?;
self.put_expression(sampler, context, true)?;
write!(self.out, ", ")?;
self.put_expression(coordinate, context, true)?;
write!(self.out, ")")?;
}
crate::Expression::ImageSample {
image,
sampler,
gather,
coordinate,
array_index,
offset,
level,
depth_ref,
clamp_to_edge,
} => { if clamp_to_edge { return Err(Error::GenericValidation( "ImageSample::clamp_to_edge should have been validated out".to_string(),
));
}
let main_op = match gather {
Some(_) => "gather",
None => "sample",
};
let comparison_op = match depth_ref {
Some(_) => "_compare",
None => "",
};
self.put_expression(image, context, false)?;
write!(self.out, ".{main_op}{comparison_op}(")?;
self.put_expression(sampler, context, true)?;
write!(self.out, ", ")?;
self.put_expression(coordinate, context, true)?; if let Some(expr) = array_index {
write!(self.out, ", ")?;
self.put_expression(expr, context, true)?;
} if let Some(dref) = depth_ref {
write!(self.out, ", ")?;
self.put_expression(dref, context, true)?;
}
return self.put_dot_product(
Reinterpreted::new(packed_type, arg),
Reinterpreted::new(packed_type, arg1.unwrap()), 4,
|writer, arg, index| { // MSL implicitly promotes these (signed or unsigned) chars to // `int` or `uint` in the multiplication, so no overflow can occur.
write!(writer.out, "{arg}[{index}]")?;
Ok(())
},
);
} else { // Fall back to a polyfill since MSL < 2.1 doesn't seem to support // bitcasting from uint to `packed_char4` or `packed_uchar4`. // See <https://github.com/gfx-rs/wgpu/pull/7574#issuecomment-2835464472>.
let conversion = match fun {
Mf::Dot4I8Packed => "int",
Mf::Dot4U8Packed => "",
_ => unreachable!(),
};
// or metal will complain that select is ambiguous
match *inner {
crate::TypeInner::Vector { size, scalar } => {
let size = common::vector_size_str(size);
let name = scalar.to_msl_name();
write!(self.out, "{name}{size}")?;
}
crate::TypeInner::Scalar(scalar) => {
let name = scalar.to_msl_name();
write!(self.out, "{name}")?;
}
_ => (),
}
write!(self.out, "(-1), ")?;
self.put_expression(arg, context, true)?;
write!(self.out, " == 0")?; if scalar.kind == crate::ScalarKind::Sint {
write!(self.out, " || ")?;
self.put_expression(arg, context, true)?;
write!(self.out, " == -1")?;
}
write!(self.out, ")")?;
}
Mf::Unpack2x16float => {
write!(self.out, "float2(as_type<half2>(")?;
self.put_expression(arg, context, false)?;
write!(self.out, "))")?;
}
Mf::Pack2x16float => {
write!(self.out, "as_type<uint>(half2(")?;
self.put_expression(arg, context, false)?;
write!(self.out, "))")?;
}
Mf::ExtractBits => { // The behavior of ExtractBits is undefined when offset + count > bit_width. We need // to first sanitize the offset and count first. If we don't do this, Apple chips // will return out-of-spec values if the extracted range is not within the bit width. // // This encodes the exact formula specified by the wgsl spec, without temporary values: // https://gpuweb.github.io/gpuweb/wgsl/#extractBits-unsigned-builtin // // w = sizeof(x) * 8 // o = min(offset, w) // tmp = w - o // c = min(count, tmp) // // bitfieldExtract(x, o, c) // // extract_bits(e, min(offset, w), min(count, w - min(offset, w))))
let scalar_bits = context.resolve_type(arg).scalar_width().unwrap() * 8;
write!(self.out, "{NAMESPACE}::extract_bits(")?;
self.put_expression(arg, context, true)?;
write!(self.out, ", {NAMESPACE}::min(")?;
self.put_expression(arg1.unwrap(), context, true)?;
write!(self.out, ", {scalar_bits}u), {NAMESPACE}::min(")?;
self.put_expression(arg2.unwrap(), context, true)?;
write!(self.out, ", {scalar_bits}u - {NAMESPACE}::min(")?;
self.put_expression(arg1.unwrap(), context, true)?;
write!(self.out, ", {scalar_bits}u)))")?;
}
Mf::InsertBits => { // The behavior of InsertBits has the same issue as ExtractBits. // // insertBits(e, newBits, min(offset, w), min(count, w - min(offset, w))))
let scalar_bits = context.resolve_type(arg).scalar_width().unwrap() * 8;
write!(
self.out, "{}_{committed}(",
super::ray::INTERSECTION_FUNCTION_NAME
)?;
self.put_expression(query, context, true)?;
write!(self.out, ")")?;
}
crate::Expression::CooperativeLoad { ref data, .. } => { if context.lang_version < (2, 3) { return Err(Error::UnsupportedCooperativeMatrix);
}
write!(self.out, "{COOPERATIVE_LOAD_FUNCTION}(")?;
write!(self.out, "&")?;
self.put_access_chain(data.pointer, context.policies.index, context)?;
write!(self.out, ", ")?;
self.put_expression(data.stride, context, true)?; // Metal's `simdgroup_load` treats its `transpose` flag as // "memory is transposed from the simdgroup_matrix's canonical // layout". On Apple GPUs that canonical layout is row-major, // so `transpose=false` loads from row-major memory. WGSL's // `coopLoadT` (row_major=true) = row-major memory, so it must // map to `transpose=false`. Hence the negation.
write!(self.out, ", {})", !data.row_major)?;
}
crate::Expression::CooperativeMultiplyAdd { a, b, c } => { if context.lang_version < (2, 3) { return Err(Error::UnsupportedCooperativeMatrix);
}
write!(self.out, "{COOPERATIVE_MULTIPLY_ADD_FUNCTION}(")?;
self.put_expression(a, context, true)?;
write!(self.out, ", ")?;
self.put_expression(b, context, true)?;
write!(self.out, ", ")?;
self.put_expression(c, context, true)?;
write!(self.out, ")")?;
}
}
Ok(())
}
/// Emits code for a binary operation, using the provided callback to emit /// the left and right operands.
fn put_binop<F>(
&mut self,
op: crate::BinaryOperator,
left: Handle<crate::Expression>,
right: Handle<crate::Expression>,
context: &ExpressionContext,
is_scoped: bool,
put_expression: &F,
) -> BackendResult
where
F: Fn(&mut Self, Handle<crate::Expression>, &ExpressionContext, bool) -> BackendResult,
{
let op_str = back::binary_operation_str(op);
if !is_scoped {
write!(self.out, "(")?;
}
// Cast packed vector if necessary // Packed vector - matrix multiplications are not supported in MSL if op == crate::BinaryOperator::Multiply
&& matches!(
context.resolve_type(right),
&crate::TypeInner::Matrix { .. }
)
{
self.put_wrapped_expression_for_packed_vec3_access(
left,
context, false,
put_expression,
)?;
} else {
put_expression(self, left, context, false)?;
}
write!(self.out, " {op_str} ")?;
// See comment above if op == crate::BinaryOperator::Multiply
&& matches!(context.resolve_type(left), &crate::TypeInner::Matrix { .. })
{
self.put_wrapped_expression_for_packed_vec3_access(
right,
context, false,
put_expression,
)?;
} else {
put_expression(self, right, context, false)?;
}
if !is_scoped {
write!(self.out, ")")?;
}
Ok(())
}
/// Used by expressions like Swizzle and Binary since they need packed_vec3's to be casted to a vec3
fn put_wrapped_expression_for_packed_vec3_access<F>(
&mut self,
expr_handle: Handle<crate::Expression>,
context: &ExpressionContext,
is_scoped: bool,
put_expression: &F,
) -> BackendResult
where
F: Fn(&mut Self, Handle<crate::Expression>, &ExpressionContext, bool) -> BackendResult,
{ if let Some(scalar) = context.get_packed_vec_kind(expr_handle) {
write!(self.out, "{}::{}3(", NAMESPACE, scalar.to_msl_name())?;
put_expression(self, expr_handle, context, is_scoped)?;
write!(self.out, ")")?;
} else {
put_expression(self, expr_handle, context, is_scoped)?;
}
Ok(())
}
/// Emits code for an expression using the provided callback, wrapping the /// result in a bitcast to the type `cast_to`.
fn put_bitcasted_expression<F>(
&mut self,
cast_to: &crate::TypeInner,
inner_expr: Handle<crate::Expression>,
context: &ExpressionContext,
put_expression: &F,
) -> BackendResult
where
F: Fn(&mut Self, &ExpressionContext, bool) -> BackendResult,
{ // For sub-32-bit types, C++ integer promotion can widen the inner // expression (e.g. `ushort + ushort` promotes to `int`), making a // direct `as_type<short>(int_expr)` invalid due to size mismatch. // We wrap with `static_cast` to truncate back before the bitcast.
let needs_truncation = match *cast_to {
crate::TypeInner::Scalar(scalar) => scalar.width < 4,
crate::TypeInner::Vector { scalar, .. } => scalar.width < 4,
_ => false,
};
if needs_truncation {
write!(self.out, "static_cast<")?; // Cast to the unsigned version of the target type to truncate
let unsigned_scalar = match *cast_to {
crate::TypeInner::Scalar(scalar) => crate::Scalar {
kind: crate::ScalarKind::Uint,
..scalar
},
crate::TypeInner::Vector { scalar, .. } => crate::Scalar {
kind: crate::ScalarKind::Uint,
..scalar
},
_ => unreachable!(),
};
match *cast_to {
crate::TypeInner::Scalar(_) => {
put_numeric_type(&mut self.out, unsigned_scalar, &[])?
}
crate::TypeInner::Vector { size, .. } => {
put_numeric_type(&mut self.out, unsigned_scalar, &[size])?
}
_ => unreachable!(),
};
write!(self.out, ">(")?;
}
// if it's packed, we must unpack it (e.g., float3(val)) before the bitcast. if let Some(scalar) = context.get_packed_vec_kind(inner_expr) {
put_numeric_type(&mut self.out, scalar, &[crate::VectorSize::Tri])?;
write!(self.out, "(")?;
put_expression(self, context, true)?;
write!(self.out, ")")?;
} else {
put_expression(self, context, true)?;
}
if needs_truncation {
write!(self.out, ")")?;
}
write!(self.out, ")")?;
Ok(())
}
/// Write a `GuardedIndex` as a Metal expression.
fn put_index(
&mut self,
index: index::GuardedIndex,
context: &ExpressionContext,
is_scoped: bool,
) -> BackendResult {
match index {
index::GuardedIndex::Expression(expr) => {
self.put_expression(expr, context, is_scoped)?
}
index::GuardedIndex::Known(value) => write!(self.out, "{value}")?,
}
Ok(())
}
/// Emit an index bounds check condition for `chain`, if required. /// /// `chain` is a subtree of `Access` and `AccessIndex` expressions, /// operating either on a pointer to a value, or on a value directly. If we cannot /// statically determine that all indexing operations in `chain` are within /// bounds, then write a conditional expression to check them dynamically, /// and return true. All accesses in the chain are checked by the generated /// expression. /// /// This assumes that the [`BoundsCheckPolicy`] for `chain` is [`ReadZeroSkipWrite`]. /// /// The text written is of the form: /// /// ```ignore /// {level}{prefix}uint(i) < 4 && uint(j) < 10 /// ``` /// /// where `{level}` and `{prefix}` are the arguments to this function. For [`Store`] /// statements, presumably these arguments start an indented `if` statement; for /// [`Load`] expressions, the caller is probably building up a ternary `?:` /// expression. In either case, what is written is not a complete syntactic structure /// in its own right, and the caller will have to finish it off if we return `true`. /// /// If no expression is written, return false. /// /// [`BoundsCheckPolicy`]: index::BoundsCheckPolicy /// [`ReadZeroSkipWrite`]: index::BoundsCheckPolicy::ReadZeroSkipWrite /// [`Store`]: crate::Statement::Store /// [`Load`]: crate::Expression::Load
fn put_bounds_checks(
&mut self,
chain: Handle<crate::Expression>,
context: &ExpressionContext,
level: back::Level,
prefix: &'static str,
) -> Result<bool, Error> {
let mut check_written = false;
// Iterate over the access chain, handling each required bounds check. for item in context.bounds_check_iter(chain) {
let BoundsCheck {
base,
index,
length,
} = item;
// Check that the index falls within bounds. Do this with a single // comparison, by casting the index to `uint` first, so that negative // indices become large positive values.
write!(self.out, "uint(")?;
self.put_index(index, context, true)?;
self.out.write_str(") < ")?;
match length {
index::IndexableLength::Known(value) => write!(self.out, "{value}")?,
index::IndexableLength::Dynamic => {
let global = context.function.originating_global(base).ok_or_else(|| {
Error::GenericValidation("Could not find originating global".into())
})?;
write!(self.out, "1 + ")?;
self.put_dynamic_array_max_index(global, context)?
}
}
}
Ok(check_written)
}
/// Write the access chain `chain`. /// /// `chain` is a subtree of [`Access`] and [`AccessIndex`] expressions, /// operating either on a pointer to a value, or on a value directly. /// /// Generate bounds checks code only if `policy` is [`Restrict`]. The /// [`ReadZeroSkipWrite`] policy requires checks before any accesses take place, so /// that must be handled in the caller. /// /// Handle the entire chain, recursing back into `put_expression` only for index /// expressions and the base expression that originates the pointer or composite value /// being accessed. This allows `put_expression` to assume that any `Access` or /// `AccessIndex` expressions it sees are the top of a chain, so it can emit /// `ReadZeroSkipWrite` checks. /// /// [`Access`]: crate::Expression::Access /// [`AccessIndex`]: crate::Expression::AccessIndex /// [`Restrict`]: crate::proc::index::BoundsCheckPolicy::Restrict /// [`ReadZeroSkipWrite`]: crate::proc::index::BoundsCheckPolicy::ReadZeroSkipWrite
fn put_access_chain(
&mut self,
chain: Handle<crate::Expression>,
policy: index::BoundsCheckPolicy,
context: &ExpressionContext,
) -> BackendResult {
match context.function.expressions[chain] {
crate::Expression::Access { base, index } => {
let mut base_ty = context.resolve_type(base);
// Look through any pointers to see what we're really indexing. if let crate::TypeInner::Pointer { base, space: _ } = *base_ty {
base_ty = &context.module.types[base].inner;
}
self.put_subscripted_access_chain(
base,
base_ty,
index::GuardedIndex::Expression(index),
policy,
context,
)?;
}
crate::Expression::AccessIndex { base, index } => {
let base_resolution = &context.info[base].ty;
let mut base_ty = base_resolution.inner_with(&context.module.types);
let mut base_ty_handle = base_resolution.handle();
// Look through any pointers to see what we're really indexing. if let crate::TypeInner::Pointer { base, space: _ } = *base_ty {
base_ty = &context.module.types[base].inner;
base_ty_handle = Some(base);
}
// Handle structs and anything else that can use `.x` syntax here, so // `put_subscripted_access_chain` won't have to handle the absurd case of // indexing a struct with an expression.
match *base_ty {
crate::TypeInner::Struct { .. } => {
let base_ty = base_ty_handle.unwrap();
self.put_access_chain(base, policy, context)?;
let name = &self.names[&NameKey::StructMember(base_ty, index)];
write!(self.out, ".{name}")?;
}
crate::TypeInner::ValuePointer { .. } | crate::TypeInner::Vector { .. } => {
self.put_access_chain(base, policy, context)?; // Prior to Metal v2.1 component access for packed vectors wasn't available // however array indexing is if context.get_packed_vec_kind(base).is_some() {
write!(self.out, "[{index}]")?;
} else {
write!(self.out, ".{}", back::COMPONENTS[index as usize])?;
}
}
_ => {
self.put_subscripted_access_chain(
base,
base_ty,
index::GuardedIndex::Known(index),
policy,
context,
)?;
}
}
}
_ => self.put_expression(chain, context, false)?,
}
Ok(())
}
/// Write a `[]`-style access of `base` by `index`. /// /// If `policy` is [`Restrict`], then generate code as needed to force all index /// values within bounds. /// /// The `base_ty` argument must be the type we are actually indexing, like [`Array`] or /// [`Vector`]. In other words, it's `base`'s type with any surrounding [`Pointer`] /// removed. Our callers often already have this handy. /// /// This only emits `[]` expressions; it doesn't handle struct member accesses or /// referencing vector components by name. /// /// [`Restrict`]: crate::proc::index::BoundsCheckPolicy::Restrict /// [`Array`]: crate::TypeInner::Array /// [`Vector`]: crate::TypeInner::Vector /// [`Pointer`]: crate::TypeInner::Pointer
fn put_subscripted_access_chain(
&mut self,
base: Handle<crate::Expression>,
base_ty: &crate::TypeInner,
index: index::GuardedIndex,
policy: index::BoundsCheckPolicy,
context: &ExpressionContext,
) -> BackendResult {
let accessing_wrapped_array = match *base_ty {
crate::TypeInner::Array {
size: crate::ArraySize::Constant(_) | crate::ArraySize::Pending(_),
..
} => true,
_ => false,
};
let accessing_wrapped_binding_array =
matches!(*base_ty, crate::TypeInner::BindingArray { .. });
if is_atomic_pointer {
write!(
self.out, "{NAMESPACE}::atomic_load_explicit({ATOMIC_REFERENCE}"
)?;
self.put_access_chain(pointer, policy, context)?;
write!(self.out, ", {NAMESPACE}::memory_order_relaxed)")?;
} else { // We don't do any dereferencing with `*` here as pointer arguments to functions // are done by `&` references and not `*` pointers. These do not need to be // dereferenced.
self.put_access_chain(pointer, policy, context)?;
}
for (index, member) in members.iter().enumerate() { if let Some(crate::Binding::BuiltIn(crate::BuiltIn::PointSize)) =
member.binding
{
has_point_size = true; if !context.pipeline_options.allow_and_force_point_size { continue;
}
}
let comma = if is_first { "" } else { "," };
is_first = false;
let name = &self.names[&NameKey::StructMember(result_ty, index as u32)]; // HACK: we are forcefully deduplicating the expression here // to convert from a wrapped struct to a raw array, e.g. // `float gl_ClipDistance1 [[clip_distance]] [1];`. if let crate::TypeInner::Array {
size: crate::ArraySize::Constant(size),
..
} = context.module.types[member.ty].inner
{
write!(self.out, "{comma} {{")?; for j in 0..size.get() { if j != 0 {
write!(self.out, ",")?;
}
write!(self.out, "{tmp}.{name}.{WRAPPED_ARRAY_FIELD}[{j}]")?;
}
write!(self.out, "}}")?;
} else {
write!(self.out, "{comma} {tmp}.{name}")?;
}
}
}
_ => {
write!(self.out, "{level}return {struct_name} {{ ")?;
self.put_expression(expr_handle, context, true)?;
}
}
if let FunctionOrigin::EntryPoint(ep_index) = context.origin {
let stage = context.module.entry_points[ep_index as usize].stage; if context.pipeline_options.allow_and_force_point_size
&& stage == crate::ShaderStage::Vertex
&& !has_point_size
{ // point size was injected and comes last
write!(self.out, ", 1.0")?;
}
}
write!(self.out, " }}")?;
}
None => {
write!(self.out, "{level}return ")?;
self.put_expression(expr_handle, context, true)?;
}
}
writeln!(self.out, ";")?;
Ok(())
}
/// Helper method used to find which expressions of a given function require baking /// /// # Notes /// This function overwrites the contents of `self.need_bake_expressions`
fn update_expressions_to_bake(
&mut self,
func: &crate::Function,
info: &valid::FunctionInfo,
context: &ExpressionContext,
) {
use crate::Expression;
self.need_bake_expressions.clear();
for (expr_handle, expr) in func.expressions.iter() { // Expressions whose reference count is above the // threshold should always be stored in temporaries.
let expr_info = &info[expr_handle];
let min_ref_count = func.expressions[expr_handle].bake_ref_count(); if min_ref_count <= expr_info.ref_count {
self.need_bake_expressions.insert(expr_handle);
} else {
match expr_info.ty { // force ray desc to be baked: it's used multiple times internally
TypeResolution::Handle(h) if Some(h) == context.module.special_types.ray_desc =>
{
self.need_bake_expressions.insert(expr_handle);
}
_ => {}
}
}
if let Expression::Math {
fun,
arg,
arg1,
arg2,
..
} = *expr
{
match fun { // WGSL's `dot` function works on any `vecN` type, but Metal's only // works on floating-point vectors, so we emit inline code for // integer vector `dot` calls. But that code uses each argument `N` // times, once for each component (see `put_dot_product`), so to // avoid duplicated evaluation, we must bake integer operands. // This applies both when using the polyfill (because of the duplicate // evaluation issue) and when we don't use the polyfill (because we // need them to be emitted before casting to packed chars -- see the // comment at the call to `put_casting_to_packed_chars`).
crate::MathFunction::Dot4U8Packed | crate::MathFunction::Dot4I8Packed => {
self.need_bake_expressions.insert(arg);
self.need_bake_expressions.insert(arg1.unwrap());
}
crate::MathFunction::FirstLeadingBit => {
self.need_bake_expressions.insert(arg);
}
crate::MathFunction::Pack4xI8
| crate::MathFunction::Pack4xU8
| crate::MathFunction::Pack4xI8Clamp
| crate::MathFunction::Pack4xU8Clamp
| crate::MathFunction::Unpack4xI8
| crate::MathFunction::Unpack4xU8 => { // On MSL < 2.1, we emit a polyfill for these functions that uses the // argument multiple times. This is no longer necessary on MSL >= 2.1. if context.lang_version < (2, 1) {
self.need_bake_expressions.insert(arg);
}
}
crate::MathFunction::ExtractBits => { // Only argument 1 is re-used.
self.need_bake_expressions.insert(arg1.unwrap());
}
crate::MathFunction::InsertBits => { // Only argument 2 is re-used.
self.need_bake_expressions.insert(arg2.unwrap());
}
crate::MathFunction::Sign => { // WGSL's `sign` function works also on signed ints, but Metal's only // works on floating points, so we emit inline code for integer `sign` // calls. But that code uses each argument 2 times (see `put_isign`), // so to avoid duplicated evaluation, we must bake the argument.
let inner = context.resolve_type(expr_handle); if inner.scalar_kind() == Some(crate::ScalarKind::Sint) {
self.need_bake_expressions.insert(arg);
}
}
_ => {}
}
}
}
}
//TODO: figure out the naming scheme that wouldn't collide with user names.
write!(self.out, " {name} = ")?;
Ok(())
}
/// Cache a clamped level of detail value, if necessary. /// /// [`ImageLoad`] accesses covered by [`BoundsCheckPolicy::Restrict`] use a /// properly clamped level of detail value both in the access itself, and /// for fetching the size of the requested MIP level, needed to clamp the /// coordinates. To avoid recomputing this clamped level of detail, we cache /// it in a temporary variable, as part of the [`Emit`] statement covering /// the [`ImageLoad`] expression. /// /// [`ImageLoad`]: crate::Expression::ImageLoad /// [`BoundsCheckPolicy::Restrict`]: index::BoundsCheckPolicy::Restrict /// [`Emit`]: crate::Statement::Emit
fn put_cache_restricted_level(
&mut self,
load: Handle<crate::Expression>,
image: Handle<crate::Expression>,
mip_level: Option<Handle<crate::Expression>>,
indent: back::Level,
context: &StatementContext,
) -> BackendResult { // Does this image access actually require (or even permit) a // level-of-detail, and does the policy require us to restrict it?
let level_of_detail = match mip_level {
Some(level) => level,
None => return Ok(()),
};
if context.expression.policies.image_load != index::BoundsCheckPolicy::Restrict
|| !context.expression.image_needs_lod(image)
{ return Ok(());
}
/// Convert the arguments of `Dot4{I, U}Packed` to `packed_(u?)char4`. /// /// Caches the results in temporary variables (whose names are derived from /// the original variable names). This caching avoids the need to redo the /// casting for each vector component when emitting the dot product.
fn put_casting_to_packed_chars(
&mut self,
fun: crate::MathFunction,
arg0: Handle<crate::Expression>,
arg1: Handle<crate::Expression>,
indent: back::Level,
context: &StatementContext<'_>,
) -> Result<(), Error> {
let packed_type = match fun {
crate::MathFunction::Dot4I8Packed => "packed_char4",
crate::MathFunction::Dot4U8Packed => "packed_uchar4",
_ => unreachable!(),
};
fn put_block(
&mut self,
level: back::Level,
statements: &[crate::Statement],
context: &StatementContext,
) -> BackendResult { // Add to the set in order to track the stack size. #[cfg(test)]
self.put_block_stack_pointers
.insert(ptr::from_ref(&level).cast());
for statement in statements {
log::trace!("statement[{}] {:?}", level.0, statement);
match *statement {
crate::Statement::Emit(ref range) => { for handle in range.clone() {
use crate::MathFunction as Mf;
match context.expression.function.expressions[handle] { // `ImageLoad` expressions covered by the `Restrict` bounds check policy // may need to cache a clamped version of their level-of-detail argument.
crate::Expression::ImageLoad {
image,
level: mip_level,
..
} => {
self.put_cache_restricted_level(
handle, image, mip_level, level, context,
)?;
}
// If we are going to write a `Dot4I8Packed` or `Dot4U8Packed` on Metal // 2.1+ then we introduce two intermediate variables that recast the two // arguments as packed (signed or unsigned) chars. The actual dot product // is implemented in `Self::put_expression`, and it uses both of these // intermediate variables multiple times. There's no danger that the // original arguments get modified between the definition of these // intermediate variables and the implementation of the actual dot // product since we require the inputs of `Dot4{I, U}Packed` to be baked.
crate::Expression::Math {
fun: fun @ (Mf::Dot4I8Packed | Mf::Dot4U8Packed),
arg,
arg1,
..
} if context.expression.lang_version >= (2, 1) => {
self.put_casting_to_packed_chars(
fun,
arg,
arg1.unwrap(),
level,
context,
)?;
}
_ => (),
}
let ptr_class = context.expression.resolve_type(handle).pointer_space();
let expr_name = if ptr_class.is_some() {
None // don't bake pointer expressions (just yet)
} elseif let Some(name) =
context.expression.function.named_expressions.get(&handle)
{ // The `crate::Function::named_expressions` table holds // expressions that should be saved in temporaries once they // are `Emit`ted. We only add them to `self.named_expressions` // when we reach the `Emit` that covers them, so that we don't // try to use their names before we've actually initialized // the temporary that holds them. // // Don't assume the names in `named_expressions` are unique, // or even valid. Use the `Namer`.
Some(self.namer.call(name))
} else { // If this expression is an index that we're going to first compare // against a limit, and then actually use as an index, then we may // want to cache it in a temporary, to avoid evaluating it twice.
let bake = if context.expression.guarded_indices.contains(handle) {
true
} else {
self.need_bake_expressions.contains(&handle)
};
if bake {
Some(Baked(handle).to_string())
} else {
None
}
};
// This backend supports `SHADER_INT64_ATOMIC_MIN_MAX` but not // `SHADER_INT64_ATOMIC_ALL_OPS`, so we can assume that if `result` is // `Some`, we are not operating on a 64-bit value, and that if we are // operating on a 64-bit value, `result` is `None`.
write!(self.out, "{level}")?;
let fun_key = if let Some(result) = result {
let res_name = Baked(result).to_string();
self.start_baking_expression(result, context, &res_name)?;
self.named_expressions.insert(result, res_name);
fun.to_msl()
} elseif context.resolve_type(value).scalar_width() == Some(8) {
fun.to_msl_64_bit()?
} else {
fun.to_msl()
};
// If the pointer we're passing to the atomic operation needs to be conditional // for `ReadZeroSkipWrite`, the condition needs to *surround* the atomic op, and // the pointer operand should be unchecked.
let policy = context.choose_bounds_check_policy(pointer);
let checked = policy == index::BoundsCheckPolicy::ReadZeroSkipWrite
&& self.put_bounds_checks(pointer, context, back::Level(0), "")?;
// If requested and successfully put bounds checks, continue the ternary expression. if checked {
write!(self.out, " ? ")?;
}
// un-emit expressions //TODO: take care of loop/continuing? for statement in statements { if let crate::Statement::Emit(ref range) = *statement { for handle in range.clone() {
self.named_expressions.shift_remove(&handle);
}
}
}
Ok(())
}
writeln!(
self.out, "// language: metal{}.{}",
options.lang_version.0, options.lang_version.1
)?;
writeln!(self.out, "#include <metal_stdlib>")?;
writeln!(self.out, "#include <simd/simd.h>")?;
writeln!(self.out)?; // Work around Metal bug where `uint` is not available by default
writeln!(self.out, "using {NAMESPACE}::uint;")?;
if module.special_types.ray_desc.is_some()
|| module.special_types.ray_intersection.is_some()
{ if options.lang_version < (2, 4) { return Err(Error::UnsupportedRayTracing);
}
}
if options
.bounds_check_policies
.contains(index::BoundsCheckPolicy::ReadZeroSkipWrite)
{
self.put_default_constructible()?;
}
writeln!(self.out)?;
{ // Make a `Vec` of all the `GlobalVariable`s that contain // runtime-sized arrays.
let globals: Vec<Handle<crate::GlobalVariable>> = module
.global_variables
.iter()
.filter(|&(_, var)| needs_array_length(var.ty, &module.types))
.map(|(handle, _)| handle)
.collect();
let mut buffer_indices = vec![]; for vbm in &pipeline_options.vertex_buffer_mappings {
buffer_indices.push(vbm.id);
}
if !globals.is_empty() || !buffer_indices.is_empty() {
writeln!(self.out, "struct _mslBufferSizes {{")?;
for global in globals {
writeln!(
self.out, "{}uint {};",
back::INDENT,
ArraySizeMember(global)
)?;
}
for idx in buffer_indices {
writeln!(self.out, "{}uint buffer_size{};", back::INDENT, idx)?;
}
/// Write the definition for the `DefaultConstructible` class. /// /// The [`ReadZeroSkipWrite`] bounds check policy requires us to be able to /// produce 'zero' values for any type, including structs, arrays, and so /// on. We could do this by emitting default constructor applications, but /// that would entail printing the name of the type, which is more trouble /// than you'd think. Instead, we just construct this magic C++14 class that /// can be converted to any type that can be default constructed, using /// template parameter inference to detect which type is needed, so we don't /// have to figure out the name. /// /// [`ReadZeroSkipWrite`]: index::BoundsCheckPolicy::ReadZeroSkipWrite
fn put_default_constructible(&mut self) -> BackendResult {
let tab = back::INDENT;
writeln!(self.out, "struct DefaultConstructible {{")?;
writeln!(self.out, "{tab}template<typename T>")?;
writeln!(self.out, "{tab}operator T() && {{")?;
writeln!(self.out, "{tab}{tab}return T {{}};")?;
writeln!(self.out, "{tab}}}")?;
writeln!(self.out, "}};")?;
Ok(())
}
if !ty.needs_alias() { continue;
}
let name = &self.names[&NameKey::Type(handle)];
match ty.inner { // Naga IR can pass around arrays by value, but Metal, following // C++, performs an array-to-pointer conversion (C++ [conv.array]) // on expressions of array type, so assigning the array by value // isn't possible. However, Metal *does* assign structs by // value. So in our Metal output, we wrap all array types in // synthetic struct types: // // struct type1 { // float inner[10] // }; // // Then we carefully include `.inner` (`WRAPPED_ARRAY_FIELD`) in // any expression that actually wants access to the array.
crate::TypeInner::Array {
base,
size,
stride: _,
} => {
let base_name = TypeContext {
handle: base,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: false,
};
match size.resolve(module.to_ctx())? {
proc::IndexableLength::Known(size) => {
writeln!(self.out, "struct {name} {{")?;
writeln!(
self.out, "{}{} {}[{}];",
back::INDENT,
base_name,
WRAPPED_ARRAY_FIELD,
size
)?;
writeln!(self.out, "}};")?;
}
proc::IndexableLength::Dynamic => {
writeln!(self.out, "typedef {base_name} {name}[1];")?;
}
}
}
crate::TypeInner::Struct {
ref members, span, ..
} => {
writeln!(self.out, "struct {name} {{")?;
let mut last_offset = 0; for (index, member) in members.iter().enumerate() { if member.offset > last_offset {
self.struct_member_pads.insert((handle, index as u32));
let pad = member.offset - last_offset;
writeln!(self.out, "{}char _pad{}[{}];", back::INDENT, index, pad)?;
}
let ty_inner = &module.types[member.ty].inner;
last_offset = member.offset + ty_inner.size(module.to_ctx());
let member_name = &self.names[&NameKey::StructMember(handle, index as u32)];
// If the member should be packed (as is the case for a misaligned vec3) issue a packed vector
match should_pack_struct_member(members, span, index, module) {
Some(scalar) => {
writeln!(
self.out, "{}{}::packed_{}3 {};",
back::INDENT, NAMESPACE,
scalar.to_msl_name(),
member_name
)?;
}
None => {
let base_name = TypeContext {
handle: member.ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: false,
};
writeln!(
self.out, "{}{} {};",
back::INDENT,
base_name,
member_name
)?;
// for 3-component vectors, add one component if let crate::TypeInner::Vector {
size: crate::VectorSize::Tri,
scalar,
} = *ty_inner
{
last_offset += scalar.width as u32;
}
}
}
} if last_offset < span {
let pad = span - last_offset;
writeln!(
self.out, "{}char _pad{}[{}];",
back::INDENT,
members.len(),
pad
)?;
}
writeln!(self.out, "}};")?;
}
_ => {
let ty_name = TypeContext {
handle,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: true,
};
writeln!(self.out, "typedef {ty_name} {name};")?;
}
}
}
// Write functions to create special types. for (type_key, struct_ty) in module.special_types.predeclared_types.iter() {
match type_key {
&crate::PredeclaredType::ModfResult { size, scalar }
| &crate::PredeclaredType::FrexpResult { size, scalar } => {
let arg_type_name_owner;
let arg_type_name = if let Some(size) = size {
arg_type_name_owner = format!( "{NAMESPACE}::{}{}", if scalar.width == 8 { "double" } else { "float" },
size as u8
);
&arg_type_name_owner
} elseif scalar.width == 8 { "double"
} else { "float"
};
let other_type_name_owner;
let (defined_func_name, called_func_name, other_type_name) = if matches!(type_key, &crate::PredeclaredType::ModfResult { .. }) {
(MODF_FUNCTION, "modf", arg_type_name)
} else {
let other_type_name = if let Some(size) = size {
other_type_name_owner = format!("int{}", size as u8);
&other_type_name_owner
} else { "int"
};
(FREXP_FUNCTION, "frexp", other_type_name)
};
let struct_name = &self.names[&NameKey::Type(*struct_ty)];
writeln!(self.out)?;
writeln!(
self.out, "{struct_name} {defined_func_name}({arg_type_name} arg) {{
{other_type_name} other;
{arg_type_name} fract = {NAMESPACE}::{called_func_name}(arg, other); return {struct_name}{{ fract, other }};
}}"
)?;
}
&crate::PredeclaredType::AtomicCompareExchangeWeakResult(scalar) => {
let arg_type_name = scalar.to_msl_name();
let called_func_name = "atomic_compare_exchange_weak_explicit";
let defined_func_name = ATOMIC_COMP_EXCH_FUNCTION;
let struct_name = &self.names[&NameKey::Type(*struct_ty)];
writeln!(self.out)?;
for address_space_name in ["device", "threadgroup"] {
writeln!(
self.out, "\ template <typename A>
{struct_name} {defined_func_name}(
{address_space_name} A *atomic_ptr,
{arg_type_name} cmp,
{arg_type_name} v
) {{ bool swapped = {NAMESPACE}::{called_func_name}(
atomic_ptr, &cmp, v,
metal::memory_order_relaxed, metal::memory_order_relaxed
); return {struct_name}{{cmp, swapped}};
}}"
)?;
}
}
}
}
Ok(())
}
/// Writes all named constants
fn write_global_constants(
&mut self,
module: &crate::Module,
mod_info: &valid::ModuleInfo,
) -> BackendResult {
let constants = module.constants.iter().filter(|&(_, c)| c.name.is_some());
for (handle, constant) in constants {
let ty_name = TypeContext {
handle: constant.ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: false,
};
let name = &self.names[&NameKey::Constant(handle)];
write!(self.out, "constant {ty_name} {name} = ")?;
self.put_const_expression(constant.init, module, mod_info, &module.global_expressions)?;
writeln!(self.out, ";")?;
}
Ok(())
}
fn put_inline_sampler_properties(
&mut self,
level: back::Level,
sampler: &sm::InlineSampler,
) -> BackendResult { for (&letter, address) in ['s', 't', 'r'].iter().zip(sampler.address.iter()) {
writeln!(
self.out, "{}{}::{}_address::{},",
level, NAMESPACE,
letter,
address.as_str(),
)?;
}
writeln!(
self.out, "{}{}::mag_filter::{},",
level, NAMESPACE,
sampler.mag_filter.as_str(),
)?;
writeln!(
self.out, "{}{}::min_filter::{},",
level, NAMESPACE,
sampler.min_filter.as_str(),
)?; if let Some(filter) = sampler.mip_filter {
writeln!(
self.out, "{}{}::mip_filter::{},",
level, NAMESPACE,
filter.as_str(),
)?;
} // avoid setting it on platforms that don't support it if sampler.border_color != sm::BorderColor::TransparentBlack {
writeln!(
self.out, "{}{}::border_color::{},",
level, NAMESPACE,
sampler.border_color.as_str(),
)?;
} //TODO: I'm not able to feed this in a way that MSL likes: //>error: use of undeclared identifier 'lod_clamp' //>error: no member named 'max_anisotropy' in namespace 'metal' iffalse { if let Some(ref lod) = sampler.lod_clamp {
writeln!(self.out, "{}lod_clamp({},{}),", level, lod.start, lod.end,)?;
} if let Some(aniso) = sampler.max_anisotropy {
writeln!(self.out, "{}max_anisotropy({}),", level, aniso.get(),)?;
}
} if sampler.compare_func != sm::CompareFunc::Never {
writeln!(
self.out, "{}{}::compare_func::{},",
level, NAMESPACE,
sampler.compare_func.as_str(),
)?;
}
writeln!(
self.out, "{}{}::coord::{}",
level, NAMESPACE,
sampler.coord.as_str()
)?;
Ok(())
}
fn write_wrapped_unary_op(
&mut self,
module: &crate::Module,
func_ctx: &back::FunctionCtx,
op: crate::UnaryOperator,
operand: Handle<crate::Expression>,
) -> BackendResult {
let operand_ty = func_ctx.resolve_type(operand, &module.types);
match op { // Negating the TYPE_MIN of a two's complement signed integer // type causes overflow, which is undefined behaviour in MSL. To // avoid this we bitcast the value to unsigned and negate it, // then bitcast back to signed. // This adheres to the WGSL spec in that the negative of the // type's minimum value should equal to the minimum value.
crate::UnaryOperator::Negate if operand_ty.scalar_kind() == Some(crate::ScalarKind::Sint) =>
{
let Some((vector_size, scalar)) = operand_ty.vector_size_and_scalar() else { return Ok(());
};
let wrapped = WrappedFunction::UnaryOp {
op,
ty: (vector_size, scalar),
}; if !self.wrapped_functions.insert(wrapped) { return Ok(());
}
writeln!(self.out, "{type_name} {NEG_FUNCTION}({type_name} val) {{")?;
let level = back::Level(1); // For sub-32-bit types, C++ integer promotion widens // `-as_type<ushort>(val)` to `int`, so we need static_cast // to truncate back before the outer as_type bitcast. if scalar.width < 4 {
writeln!(
self.out, "{level}return as_type<{type_name}>(static_cast<{unsigned_type_name}>(-as_type<{unsigned_type_name}>(val)));"
)?;
} else {
writeln!(
self.out, "{level}return as_type<{type_name}>(-as_type<{unsigned_type_name}>(val));"
)?;
}
writeln!(self.out, "}}")?;
writeln!(self.out)?;
}
_ => {}
}
Ok(())
}
fn write_wrapped_binary_op(
&mut self,
module: &crate::Module,
func_ctx: &back::FunctionCtx,
expr: Handle<crate::Expression>,
op: crate::BinaryOperator,
left: Handle<crate::Expression>,
right: Handle<crate::Expression>,
) -> BackendResult {
let expr_ty = func_ctx.resolve_type(expr, &module.types);
let left_ty = func_ctx.resolve_type(left, &module.types);
let right_ty = func_ctx.resolve_type(right, &module.types);
match (op, expr_ty.scalar_kind()) { // Signed integer division of TYPE_MIN / -1, or signed or // unsigned division by zero, gives an unspecified value in MSL. // We override the divisor to 1 in these cases. // This adheres to the WGSL spec in that: // * TYPE_MIN / -1 == TYPE_MIN // * x / 0 == x
(
crate::BinaryOperator::Divide,
Some(crate::ScalarKind::Sint | crate::ScalarKind::Uint),
) => {
let Some(left_wrapped_ty) = left_ty.vector_size_and_scalar() else { return Ok(());
};
let Some(right_wrapped_ty) = right_ty.vector_size_and_scalar() else { return Ok(());
};
let wrapped = WrappedFunction::BinaryOp {
op,
left_ty: left_wrapped_ty,
right_ty: right_wrapped_ty,
}; if !self.wrapped_functions.insert(wrapped) { return Ok(());
}
let Some((vector_size, scalar)) = expr_ty.vector_size_and_scalar() else { return Ok(());
};
let mut type_name = String::new();
match vector_size {
None => put_numeric_type(&mut type_name, scalar, &[])?,
Some(size) => put_numeric_type(&mut type_name, scalar, &[size])?,
};
writeln!(
self.out, "{type_name} {DIV_FUNCTION}({type_name} lhs, {type_name} rhs) {{"
)?;
let level = back::Level(1);
// Sub-32-bit types need typed literal wrappers (e.g. `short(1)`)
// to avoid ambiguous metal::select overloads. For >= 32-bit,
// bare literals like `1`, `-1`, `0` are unambiguous.
let (lp, rp) = if scalar.width < 4 {
(format!("{type_name}("), ")".to_string())
} else {
(String::new(), String::new())
};
match scalar.kind {
crate::ScalarKind::Sint => {
let min_val = match scalar.width { 2 => crate::Literal::I16(i16::MIN), 4 => crate::Literal::I32(i32::MIN), 8 => crate::Literal::I64(i64::MIN),
_ => {
return Err(Error::GenericValidation(format!(
"Unexpected width for scalar {scalar:?}"
)));
}
};
write!(
self.out,
"{level}return lhs / metal::select(rhs, {lp}1{rp}, (lhs == "
)?;
self.put_literal(min_val)?;
writeln!(self.out, " & rhs == {lp}-1{rp}) | (rhs == {lp}0{rp}));")?
}
crate::ScalarKind::Uint => {
let suffix = if scalar.width < 4 { "" } else { "u" };
writeln!(
self.out,
"{level}return lhs / metal::select(rhs, {lp}1{suffix}{rp}, rhs == {lp}0{suffix}{rp});"
)?
}
_ => unreachable!(),
}
writeln!(self.out, "}}")?;
writeln!(self.out)?;
}
// Integer modulo where one or both operands are negative, or the
// divisor is zero, is undefined behaviour in MSL. To avoid this
// we use the following equation:
//
// dividend - (dividend / divisor) * divisor
//
// overriding the divisor to 1 if either it is 0, or it is -1
// and the dividend is TYPE_MIN.
//
// This adheres to the WGSL spec in that:
// * TYPE_MIN % -1 == 0
// * x % 0 == 0
(
crate::BinaryOperator::Modulo,
Some(crate::ScalarKind::Sint | crate::ScalarKind::Uint),
) => {
let Some(left_wrapped_ty) = left_ty.vector_size_and_scalar() else {
return Ok(());
};
let Some((right_vector_size, right_scalar)) = right_ty.vector_size_and_scalar()
else {
return Ok(());
};
let wrapped = WrappedFunction::BinaryOp {
op,
left_ty: left_wrapped_ty,
right_ty: (right_vector_size, right_scalar),
};
if !self.wrapped_functions.insert(wrapped) {
return Ok(());
}
/// Build the mangled helper name for integer vector dot products.
///
/// `scalar` must be a concrete integer scalar type.
///
/// Result format: `{DOT_FUNCTION_PREFIX}_{type}{N}` (e.g., `naga_dot_int3`).
fn get_dot_wrapper_function_helper_name(
&self,
scalar: crate::Scalar,
size: crate::VectorSize,
) -> String {
// Check for consistency with [`super::keywords::RESERVED_SET`]
debug_assert!(concrete_int_scalars().any(|s| s == scalar));
let type_name = scalar.to_msl_name();
let size_suffix = common::vector_size_str(size);
format!("{DOT_FUNCTION_PREFIX}_{type_name}{size_suffix}")
}
#[allow(clippy::too_many_arguments)]
fn write_wrapped_math_function(
&mut self,
module: &crate::Module,
func_ctx: &back::FunctionCtx,
fun: crate::MathFunction,
arg: Handle<crate::Expression>,
_arg1: Option<Handle<crate::Expression>>,
_arg2: Option<Handle<crate::Expression>>,
_arg3: Option<Handle<crate::Expression>>,
) -> BackendResult {
let arg_ty = func_ctx.resolve_type(arg, &module.types);
match fun {
// Taking the absolute value of the TYPE_MIN of a two's
// complement signed integer type causes overflow, which is
// undefined behaviour in MSL. To avoid this, when the value is
// negative we bitcast the value to unsigned and negate it, then
// bitcast back to signed.
// This adheres to the WGSL spec in that the absolute of the
// type's minimum value should equal to the minimum value.
crate::MathFunction::Abs if arg_ty.scalar_kind() == Some(crate::ScalarKind::Sint) => {
let Some((vector_size, scalar)) = arg_ty.vector_size_and_scalar() else {
return Ok(());
};
let wrapped = WrappedFunction::Math {
fun,
arg_ty: (vector_size, scalar),
};
if !self.wrapped_functions.insert(wrapped) {
return Ok(());
}
writeln!(self.out, "{type_name} {ABS_FUNCTION}({type_name} val) {{")?;
let level = back::Level(1);
let zero = if scalar.width < 4 {
format!("{type_name}(0)")
} else {
"0".to_string()
};
let neg_expr = if scalar.width < 4 {
format!(
"static_cast<{unsigned_type_name}>(-as_type<{unsigned_type_name}>(val))"
)
} else {
format!("-as_type<{unsigned_type_name}>(val)")
};
writeln!(self.out, "{level}return metal::select(as_type<{type_name}>({neg_expr}), val, val >= {zero});")?;
writeln!(self.out, "}}")?;
writeln!(self.out)?;
}
crate::MathFunction::Dot => match *arg_ty {
crate::TypeInner::Vector { size, scalar }
if matches!(
scalar.kind,
crate::ScalarKind::Sint | crate::ScalarKind::Uint
) =>
{
// De-duplicate per (fun, arg type) like other wrapped math functions
let wrapped = WrappedFunction::Math {
fun,
arg_ty: (Some(size), scalar),
};
if !self.wrapped_functions.insert(wrapped) {
return Ok(());
}
let mut vec_ty = String::new();
put_numeric_type(&mut vec_ty, scalar, &[size])?;
let mut ret_ty = String::new();
put_numeric_type(&mut ret_ty, scalar, &[])?;
let fun_name = self.get_dot_wrapper_function_helper_name(scalar, size);
// Emit function signature and body using put_dot_product for the expression
writeln!(self.out, "{ret_ty} {fun_name}({vec_ty} a, {vec_ty} b) {{")?;
let level = back::Level(1);
write!(self.out, "{level}return ")?;
self.put_dot_product("a", "b", size as usize, |writer, name, index| {
write!(writer.out, "{name}.{}", back::COMPONENTS[index])?;
Ok(())
})?;
writeln!(self.out, ";")?;
writeln!(self.out, "}}")?;
writeln!(self.out)?;
}
_ => {}
},
_ => {}
}
Ok(())
}
fn write_wrapped_cast(
&mut self,
module: &crate::Module,
func_ctx: &back::FunctionCtx,
expr: Handle<crate::Expression>,
kind: crate::ScalarKind,
convert: Option<crate::Bytes>,
) -> BackendResult {
// Avoid undefined behaviour when casting from a float to integer
// when the value is out of range for the target type. Additionally
// ensure we clamp to the correct value as per the WGSL spec.
//
// https://www.w3.org/TR/WGSL/#floating-point-conversion:
// * If X is exactly representable in the target type T, then the
// result is that value.
// * Otherwise, the result is the value in T closest to
// truncate(X) and also exactly representable in the original
// floating point type.
let src_ty = func_ctx.resolve_type(expr, &module.types);
let Some(width) = convert else {
return Ok(());
};
let Some((vector_size, src_scalar)) = src_ty.vector_size_and_scalar() else {
return Ok(());
};
let dst_scalar = crate::Scalar { kind, width };
if src_scalar.kind != crate::ScalarKind::Float
|| (dst_scalar.kind != crate::ScalarKind::Sint
&& dst_scalar.kind != crate::ScalarKind::Uint)
{
return Ok(());
}
let wrapped = WrappedFunction::Cast {
src_scalar,
vector_size,
dst_scalar,
};
if !self.wrapped_functions.insert(wrapped) {
return Ok(());
}
let (min, max) = proc::min_max_float_representable_by(src_scalar, dst_scalar);
/// Helper function used by [`Self::write_wrapped_image_load`] and
/// [`Self::write_wrapped_image_sample`] to write the shared YUV to RGB
/// conversion code for external textures. Expects the preceding code to
/// declare the Y component as a `float` variable of name `y`, the UV
/// components as a `float2` variable of name `uv`, and the external
/// texture params as a variable of name `params`. The emitted code will
/// return the result.
fn write_convert_yuv_to_rgb_and_return(
&mut self,
level: back::Level,
y: &str,
uv: &str,
params: &str,
) -> BackendResult {
let l1 = level;
let l2 = l1.next();
// Convert from YUV to non-linear RGB in the source color space.
writeln!(
self.out,
"{l1}float3 srcGammaRgb = ({params}.yuv_conversion_matrix * float4({y}, {uv}, 1.0)).rgb;"
)?;
// Apply the inverse of the source transfer function to convert to
// linear RGB in the source color space.
writeln!(self.out, "{l1}float3 srcLinearRgb = {NAMESPACE}::select(")?;
writeln!(self.out, "{l2}{NAMESPACE}::pow((srcGammaRgb + {params}.src_tf.a - 1.0) / {params}.src_tf.a, {params}.src_tf.g),")?;
writeln!(self.out, "{l2}srcGammaRgb / {params}.src_tf.k,")?;
writeln!(
self.out,
"{l2}srcGammaRgb < {params}.src_tf.k * {params}.src_tf.b);"
)?;
// Multiply by the gamut conversion matrix to convert to linear RGB in
// the destination color space.
writeln!(
self.out,
"{l1}float3 dstLinearRgb = {params}.gamut_conversion_matrix * srcLinearRgb;"
)?;
// Finally, apply the dest transfer function to convert to non-linear
// RGB in the destination color space, and return the result.
writeln!(self.out, "{l1}float3 dstGammaRgb = {NAMESPACE}::select(")?;
writeln!(self.out, "{l2}{params}.dst_tf.a * {NAMESPACE}::pow(dstLinearRgb, 1.0 / {params}.dst_tf.g) - ({params}.dst_tf.a - 1),")?;
writeln!(self.out, "{l2}{params}.dst_tf.k * dstLinearRgb,")?;
writeln!(self.out, "{l2}dstLinearRgb < {params}.dst_tf.b);")?;
// Chroma planes may be subsampled so we must scale the coords accordingly.
writeln!(
self.out,
"{l2}uint2 plane1_size = uint2(tex.plane1.get_width(), tex.plane1.get_height());"
)?;
writeln!(self.out, "{l2}uint2 plane1_coords = uint2({NAMESPACE}::floor(float2(plane0_coords) * float2(plane1_size) / float2(plane0_size)));")?;
// For multi-plane, read the Y value from plane 0
writeln!(self.out, "{l2}float y = tex.plane0.read(plane0_coords).x;")?;
writeln!(self.out, "{l2}float2 uv;")?;
writeln!(self.out, "{l2}if (tex.params.num_planes == 2u) {{")?;
// For 2 planes, read UV from interleaved plane 1
writeln!(self.out, "{l3}uv = tex.plane1.read(plane1_coords).xy;")?;
writeln!(self.out, "{l2}}} else {{")?;
// For 3 planes, read U and V from planes 1 and 2 respectively
writeln!(
self.out,
"{l2}uint2 plane2_size = uint2(tex.plane2.get_width(), tex.plane2.get_height());"
)?;
writeln!(self.out, "{l2}uint2 plane2_coords = uint2({NAMESPACE}::floor(float2(plane0_coords) * float2(plane2_size) / float2(plane0_size)));")?;
writeln!(
self.out,
"{l3}uv = float2(tex.plane1.read(plane1_coords).x, tex.plane2.read(plane2_coords).x);"
)?;
writeln!(self.out, "{l2}}}")?;
#[allow(clippy::too_many_arguments)]
fn write_wrapped_image_sample(
&mut self,
module: &crate::Module,
func_ctx: &back::FunctionCtx,
image: Handle<crate::Expression>,
_sampler: Handle<crate::Expression>,
_gather: Option<crate::SwizzleComponent>,
_coordinate: Handle<crate::Expression>,
_array_index: Option<Handle<crate::Expression>>,
_offset: Option<Handle<crate::Expression>>,
_level: crate::SampleLevel,
_depth_ref: Option<Handle<crate::Expression>>,
clamp_to_edge: bool,
) -> BackendResult {
// We currently only need to wrap textureSampleBaseClampToEdge, for
// both sampled and external textures.
if !clamp_to_edge {
return Ok(());
}
let class = match *func_ctx.resolve_type(image, &module.types) {
crate::TypeInner::Image { class, .. } => class,
_ => unreachable!(),
};
let wrapped = WrappedFunction::ImageSample {
class,
clamp_to_edge: true,
};
if !self.wrapped_functions.insert(wrapped) {
return Ok(());
}
match class {
crate::ImageClass::External => {
writeln!(self.out, "float4 {IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION}({EXTERNAL_TEXTURE_WRAPPER_STRUCT} tex, {NAMESPACE}::sampler samp, float2 coords) {{")?;
let l1 = back::Level(1);
let l2 = l1.next();
let l3 = l2.next();
writeln!(self.out, "{l1}uint2 plane0_size = uint2(tex.plane0.get_width(), tex.plane0.get_height());")?;
writeln!(
self.out,
"{l1}coords = tex.params.sample_transform * float3(coords, 1.0);"
)?;
// Calculate the sample bounds. The purported size of the texture
// (params.size) is irrelevant here as we are dealing with normalized
// coordinates. Usually we would clamp to (0,0)..(1,1). However, we must
// apply the sample transformation to that, also bearing in mind that it
// may contain a flip on either axis. We calculate and adjust for the
// half-texel separately for each plane as it depends on the actual
// texture size which may vary between planes.
writeln!(
self.out,
"{l1}float2 bounds_min = tex.params.sample_transform * float3(0.0, 0.0, 1.0);"
)?;
writeln!(
self.out,
"{l1}float2 bounds_max = tex.params.sample_transform * float3(1.0, 1.0, 1.0);"
)?;
writeln!(self.out, "{l1}float4 bounds = float4({NAMESPACE}::min(bounds_min, bounds_max), {NAMESPACE}::max(bounds_min, bounds_max));")?;
writeln!(
self.out,
"{l1}float2 plane0_half_texel = float2(0.5, 0.5) / float2(plane0_size);"
)?;
writeln!(
self.out,
"{l1}float2 plane0_coords = {NAMESPACE}::clamp(coords, bounds.xy + plane0_half_texel, bounds.zw - plane0_half_texel);"
)?;
writeln!(self.out, "{l1}if (tex.params.num_planes == 1u) {{")?;
// For single plane, simply sample from plane0
writeln!(
self.out,
"{l2}return tex.plane0.sample(samp, plane0_coords, {NAMESPACE}::level(0.0f));"
)?;
writeln!(self.out, "{l1}}} else {{")?;
writeln!(self.out, "{l2}uint2 plane1_size = uint2(tex.plane1.get_width(), tex.plane1.get_height());")?;
writeln!(
self.out,
"{l2}float2 plane1_half_texel = float2(0.5, 0.5) / float2(plane1_size);"
)?;
writeln!(
self.out,
"{l2}float2 plane1_coords = {NAMESPACE}::clamp(coords, bounds.xy + plane1_half_texel, bounds.zw - plane1_half_texel);"
)?;
// For multi-plane, sample the Y value from plane 0
writeln!(
self.out,
"{l2}float y = tex.plane0.sample(samp, plane0_coords, {NAMESPACE}::level(0.0f)).r;"
)?;
writeln!(self.out, "{l2}float2 uv = float2(0.0, 0.0);")?;
writeln!(self.out, "{l2}if (tex.params.num_planes == 2u) {{")?;
// For 2 planes, sample UV from interleaved plane 1
writeln!(
self.out,
"{l3}uv = tex.plane1.sample(samp, plane1_coords, {NAMESPACE}::level(0.0f)).xy;"
)?;
writeln!(self.out, "{l2}}} else {{")?;
// For 3 planes, sample U and V from planes 1 and 2 respectively
writeln!(self.out, "{l3}uint2 plane2_size = uint2(tex.plane2.get_width(), tex.plane2.get_height());")?;
writeln!(
self.out,
"{l3}float2 plane2_half_texel = float2(0.5, 0.5) / float2(plane2_size);"
)?;
writeln!(
self.out,
"{l3}float2 plane2_coords = {NAMESPACE}::clamp(coords, bounds.xy + plane2_half_texel, bounds.zw - plane1_half_texel);"
)?;
writeln!(self.out, "{l3}uv.x = tex.plane1.sample(samp, plane1_coords, {NAMESPACE}::level(0.0f)).x;")?;
writeln!(self.out, "{l3}uv.y = tex.plane2.sample(samp, plane2_coords, {NAMESPACE}::level(0.0f)).x;")?;
writeln!(self.out, "{l2}}}")?;
// Returns the array of mapped entry point names.
fn write_functions(
&mut self,
module: &crate::Module,
mod_info: &valid::ModuleInfo,
options: &Options,
pipeline_options: &PipelineOptions,
) -> Result<TranslationInfo, Error> {
use back::msl::VertexFormat;
// Define structs to hold resolved/generated data for vertex buffers and
// their attributes.
struct AttributeMappingResolved {
ty_name: String,
dimension: Option<crate::VectorSize>,
scalar: crate::Scalar,
name: String,
}
let mut am_resolved = FastHashMap::<u32, AttributeMappingResolved>::default();
// Define a struct to hold a named reference to a byte-unpacking function.
struct UnpackingFunction {
name: String,
byte_count: u32,
dimension: Option<crate::VectorSize>,
scalar: crate::Scalar,
}
let mut unpacking_functions = FastHashMap::<VertexFormat, UnpackingFunction>::default();
// Check if we are attempting vertex pulling. If we are, generate some
// names we'll need, and iterate the vertex buffer mappings to output
// all the conversion functions we'll need to unpack the attribute data.
// We can re-use these names for all entry points that need them, since
// those entry points also use self.namer.
let mut needs_vertex_id = false;
let v_id = self.namer.call("v_id");
let mut needs_instance_id = false;
let i_id = self.namer.call("i_id");
if pipeline_options.vertex_pulling_transform {
for vbm in &pipeline_options.vertex_buffer_mappings {
let buffer_id = vbm.id;
let buffer_stride = vbm.stride;
let buffer_ty = self.namer.call(format!("vb_{buffer_id}_type").as_str());
let buffer_param = self.namer.call(format!("vb_{buffer_id}_in").as_str());
let buffer_elem = self.namer.call(format!("vb_{buffer_id}_elem").as_str());
let fun_info = &mod_info[fun_handle];
pass_through_globals.clear();
let mut needs_buffer_sizes = false;
for (handle, var) in module.global_variables.iter() {
if !fun_info[handle].is_empty() {
if var.space.needs_pass_through() {
pass_through_globals.push(handle);
}
needs_buffer_sizes |= needs_array_length(var.ty, &module.types);
}
}
let ep_range = get_entry_points(module, pipeline_options.entry_point.as_ref())
.map_err(|(stage, name)| Error::EntryPointNotFound(stage, name))?;
let mut info = TranslationInfo {
entry_point_names: Vec::with_capacity(ep_range.len()),
};
for ep_index in ep_range {
let ep = &module.entry_points[ep_index];
let fun = &ep.function;
let fun_info = mod_info.get_entry_point(ep_index);
let mut ep_error = None;
// For vertex_id and instance_id arguments, presume that we'll
// use our generated names, but switch to the name of an
// existing @builtin param, if we find one.
let mut v_existing_id = None;
let mut i_existing_id = None;
log::trace!(
"entry point {:?}, index {:?}",
fun.name.as_deref().unwrap_or("(anonymous)"),
ep_index
);
let ctx = back::FunctionCtx {
ty: back::FunctionType::EntryPoint(ep_index as u16),
info: fun_info,
expressions: &fun.expressions,
named_expressions: &fun.named_expressions,
};
// Should this entry point be modified to do vertex pulling?
let do_vertex_pulling = can_vertex_pull
&& pipeline_options.vertex_pulling_transform
&& !pipeline_options.vertex_buffer_mappings.is_empty();
// Is any global variable used by this entry point dynamically sized?
let needs_buffer_sizes = do_vertex_pulling
|| module
.global_variables
.iter()
.filter(|&(handle, _)| !fun_info[handle].is_empty())
.any(|(_, var)| needs_array_length(var.ty, &module.types));
// skip this entry point if any global bindings are missing,
// or their types are incompatible.
if !options.fake_missing_bindings {
for (var_handle, var) in module.global_variables.iter() {
if fun_info[var_handle].is_empty() {
continue;
}
match var.space {
crate::AddressSpace::Uniform
| crate::AddressSpace::Storage { .. }
| crate::AddressSpace::Handle => {
let br = match var.binding {
Some(ref br) => br,
None => {
let var_name = var.name.clone().unwrap_or_default();
ep_error =
Some(super::EntryPointError::MissingBinding(var_name));
break;
}
};
let target = options.get_resource_binding_target(ep, br);
let good = match target {
Some(target) => {
// We intentionally don't dereference binding_arrays here,
// so that binding arrays fall to the buffer location.
if let Some(err) = ep_error {
info.entry_point_names.push(Err(err));
continue;
}
let fun_name = self.names[&NameKey::EntryPoint(ep_index as _)].clone();
info.entry_point_names.push(Ok(fun_name.clone()));
writeln!(self.out)?;
// Since `Namer.reset` wasn't expecting struct members to be
// suddenly injected into another namespace like this,
// `self.names` doesn't keep them distinct from other variables.
// Generate fresh names for these arguments, and remember the
// mapping.
let mut flattened_member_names = FastHashMap::default();
// Varyings' members get their own namespace
let mut varyings_namer = proc::Namer::default();
let mut empty_names = FastHashMap::default(); // Create a throwaway map
varyings_namer.reset(
module,
&super::keywords::RESERVED_SET,
proc::KeywordSet::empty(),
proc::CaseInsensitiveKeywordSet::empty(),
&[CLAMPED_LOD_LOAD_PREFIX],
&mut empty_names,
);
// List all the Naga `EntryPoint`'s `Function`'s arguments,
// flattening structs into their members. In Metal, we will pass
// each of these values to the entry point as a separate argument—
// except for the varyings, handled next.
let mut flattened_arguments = Vec::new();
for (arg_index, arg) in fun.arguments.iter().enumerate() {
match module.types[arg.ty].inner {
crate::TypeInner::Struct { ref members, .. } => {
for (member_index, member) in members.iter().enumerate() {
let member_index = member_index as u32;
flattened_arguments.push((
NameKey::StructMember(arg.ty, member_index),
member.ty,
member.binding.as_ref(),
));
let name_key = NameKey::StructMember(arg.ty, member_index);
let name = match member.binding {
Some(crate::Binding::Location { .. }) => {
if do_vertex_pulling {
self.namer.call(&self.names[&name_key])
} else {
varyings_namer.call(&self.names[&name_key])
}
}
_ => self.namer.call(&self.names[&name_key]),
};
flattened_member_names.insert(name_key, name);
}
}
_ => flattened_arguments.push((
NameKey::EntryPointArgument(ep_index as _, arg_index as u32),
arg.ty,
arg.binding.as_ref(),
)),
}
}
// Identify the varyings among the argument values, and maybe emit
// a struct type named `<fun>Input` to hold them. If we are doing
// vertex pulling, we instead update our attribute mapping to
// note the types, names, and zero values of the attributes.
let stage_in_name = self.namer.call(&format!("{fun_name}Input"));
let varyings_member_name = self.namer.call("varyings");
let mut has_varyings = false;
if !flattened_arguments.is_empty() {
if !do_vertex_pulling {
writeln!(self.out, "struct {stage_in_name} {{")?;
}
for &(ref name_key, ty, binding) in flattened_arguments.iter() {
let Some(binding) = binding else {
continue;
};
let name = match *name_key {
NameKey::StructMember(..) => &flattened_member_names[name_key],
_ => &self.names[name_key],
};
let ty_name = TypeContext {
handle: ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: false,
};
let resolved = options.resolve_local_binding(binding, in_mode)?;
let location = match *binding {
crate::Binding::Location { location, .. } => Some(location),
crate::Binding::BuiltIn(crate::BuiltIn::Barycentric { .. }) => None,
crate::Binding::BuiltIn(_) => continue,
};
if do_vertex_pulling {
let Some(location) = location else {
continue;
};
// Update our attribute mapping.
am_resolved.insert(
location,
AttributeMappingResolved {
ty_name: ty_name.to_string(),
dimension: ty_name.vector_size(),
scalar: ty_name.scalar().unwrap(),
name: name.to_string(),
},
);
} else {
has_varyings = true;
if let super::ResolvedBinding::User {
prefix,
index,
interpolation: Some(super::ResolvedInterpolation::PerVertex),
} = resolved
{
if options.lang_version < (4, 0) {
return Err(Error::PerVertexNotSupported);
}
write!(
self.out,
"{}{NAMESPACE}::vertex_value<{}> {name} [[user({prefix}{index})]]",
back::INDENT,
ty_name.unwrap_array()
)?;
} else {
write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?;
resolved.try_fmt(&mut self.out)?;
}
writeln!(self.out, ";")?;
}
}
if !do_vertex_pulling {
writeln!(self.out, "}};")?;
}
}
// Define a struct type named for the return value, if any, named
// `<fun>Output`.
let stage_out_name = self.namer.call(&format!("{fun_name}Output"));
let result_member_name = self.namer.call("member");
let result_type_name = match fun.result {
Some(ref result) if ep.stage != crate::ShaderStage::Task => {
let mut result_members = Vec::new();
if let crate::TypeInner::Struct { ref members, .. } =
module.types[result.ty].inner
{
for (member_index, member) in members.iter().enumerate() {
result_members.push((
&self.names[&NameKey::StructMember(result.ty, member_index as u32)],
member.ty,
member.binding.as_ref(),
));
}
} else {
result_members.push((
&result_member_name,
result.ty,
result.binding.as_ref(),
));
}
writeln!(self.out, "struct {stage_out_name} {{")?;
let mut has_point_size = false;
for (name, ty, binding) in result_members {
let ty_name = TypeContext {
handle: ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: true,
};
let binding = binding.ok_or_else(|| {
Error::GenericValidation("Expected binding, got None".into())
})?;
if let crate::Binding::BuiltIn(crate::BuiltIn::PointSize) = *binding {
has_point_size = true;
if !pipeline_options.allow_and_force_point_size {
continue;
}
}
let array_len = match module.types[ty].inner {
crate::TypeInner::Array {
size: crate::ArraySize::Constant(size),
..
} => Some(size),
_ => None,
};
let resolved = options.resolve_local_binding(binding, out_mode)?;
write!(self.out, "{}{} {}", back::INDENT, ty_name, name)?;
resolved.try_fmt(&mut self.out)?;
if let Some(array_len) = array_len {
write!(self.out, " [{array_len}]")?;
}
writeln!(self.out, ";")?;
}
if pipeline_options.allow_and_force_point_size
&& ep.stage == crate::ShaderStage::Vertex
&& !has_point_size
{
// inject the point size output last
writeln!(
self.out,
"{}float _point_size [[point_size]];",
back::INDENT
)?;
}
writeln!(self.out, "}};")?;
&stage_out_name
}
Some(ref result) if ep.stage == crate::ShaderStage::Task => {
assert_eq!(
module.types[result.ty].inner,
crate::TypeInner::Vector {
size: crate::VectorSize::Tri,
scalar: crate::Scalar::U32
}
);
"metal::uint3"
}
_ => "void",
};
let out_mesh_info = if let Some(ref mesh_info) = ep.mesh_info {
Some(self.write_mesh_output_types(
mesh_info,
&fun_name,
module,
pipeline_options.allow_and_force_point_size,
options,
)?)
} else {
None
};
// If we're doing a vertex pulling transform, define the buffer
// structure types.
if do_vertex_pulling {
for vbm in &vbm_resolved {
let buffer_stride = vbm.stride;
let buffer_ty = &vbm.ty_name;
// Define a structure of bytes of the appropriate size.
// When we access the attributes, we'll be unpacking these
// bytes at some offset.
writeln!(
self.out,
"struct {buffer_ty} {{ metal::uchar data[{buffer_stride}]; }};"
)?;
}
}
let is_wrapped = matches!(
ep.stage,
crate::ShaderStage::Task | crate::ShaderStage::Mesh
);
let fun_name = fun_name.clone();
let nested_fun_name = if is_wrapped {
self.namer.call(&format!("_{fun_name}"))
} else {
fun_name.clone()
};
// Write the entry point function's name, and begin its argument list.
if let Some(em_str) = em_str {
write!(self.out, "{em_str} ")?;
}
writeln!(self.out, "{result_type_name} {nested_fun_name}(")?;
let mut args = Vec::new();
// If we have produced a struct holding the `EntryPoint`'s
// `Function`'s arguments' varyings, pass that struct first.
if has_varyings {
args.push(EntryPointArgument {
ty_name: stage_in_name,
name: varyings_member_name.clone(),
binding: " [[stage_in]]".to_string(),
init: None,
});
}
let mut local_invocation_index = None;
// Then pass the remaining arguments not included in the varyings
// struct.
for &(ref name_key, ty, binding) in flattened_arguments.iter() {
let binding = match binding {
Some(&crate::Binding::BuiltIn(crate::BuiltIn::Barycentric { .. })) => continue,
Some(binding @ &crate::Binding::BuiltIn { .. }) => binding,
_ => continue,
};
let name = match *name_key {
NameKey::StructMember(..) => &flattened_member_names[name_key],
_ => &self.names[name_key],
};
if binding == &crate::Binding::BuiltIn(crate::BuiltIn::LocalInvocationIndex) {
local_invocation_index = Some(name_key);
}
// Those global variables used by this entry point and its callees
// get passed as arguments. `Private` globals are an exception, they
// don't outlive this invocation, so we declare them below as locals
// within the entry point.
for (handle, var) in module.global_variables.iter() {
let usage = fun_info[handle];
if usage.is_empty() || var.space == crate::AddressSpace::Private {
continue;
}
if options.lang_version < (1, 2) {
match var.space {
// This restriction is not documented in the MSL spec
// but validation will fail if it is not upheld.
//
// We infer the required version from the "Function
// Buffer Read-Writes" section of [what's new], where
// the feature sets listed correspond with the ones
// supporting MSL 1.2.
//
// [what's new]: https://developer.apple.com/library/archive/documentation/Miscellaneous/Conceptual/MetalProgrammingGuide/WhatsNewiniOS10tvOS10andOSX1012/WhatsNewiniOS10tvOS10andOSX1012.html
crate::AddressSpace::Storage { access }
if access.contains(crate::StorageAccess::STORE)
&& ep.stage == crate::ShaderStage::Fragment =>
{
return Err(Error::UnsupportedWritableStorageBuffer)
}
crate::AddressSpace::Handle => {
match module.types[var.ty].inner {
crate::TypeInner::Image {
class: crate::ImageClass::Storage { access, .. },
..
} => {
// This restriction is not documented in the MSL spec
// but validation will fail if it is not upheld.
//
// We infer the required version from the "Function
// Texture Read-Writes" section of [what's new], where
// the feature sets listed correspond with the ones
// supporting MSL 1.2.
//
// [what's new]: https://developer.apple.com/library/archive/documentation/Miscellaneous/Conceptual/MetalProgrammingGuide/WhatsNewiniOS10tvOS10andOSX1012/WhatsNewiniOS10tvOS10andOSX1012.html
if access.contains(crate::StorageAccess::STORE)
&& (ep.stage == crate::ShaderStage::Vertex
|| ep.stage == crate::ShaderStage::Fragment)
{
return Err(Error::UnsupportedWritableStorageTexture(
ep.stage,
));
}
// the resolves have already been checked for `!fake_missing_bindings` case
let resolved = match var.space {
crate::AddressSpace::Immediate => options.resolve_immediates(ep).ok(),
crate::AddressSpace::WorkGroup => None,
crate::AddressSpace::TaskPayload => Some(back::msl::ResolvedBinding::Payload),
_ => options
.resolve_resource_binding(ep, var.binding.as_ref().unwrap())
.ok(),
};
if let Some(ref resolved) = resolved {
// Inline samplers are be defined in the EP body
if resolved.as_inline_sampler(options).is_some() {
continue;
}
}
match module.types[var.ty].inner {
crate::TypeInner::Image {
class: crate::ImageClass::External,
..
} => {
// External texture global variables get lowered to 3 textures
// and a constant buffer. We must emit a separate argument for
// each of these.
let target = match resolved {
Some(back::msl::ResolvedBinding::Resource(target)) => {
target.external_texture
}
_ => None,
};
for i in 0..3 {
let plane_name = &self.names[&NameKey::ExternalTextureGlobalVariable(
handle,
ExternalTextureNameKey::Plane(i),
)];
let ty_name = format!(
"{NAMESPACE}::texture2d<float, {NAMESPACE}::access::sample>"
);
let name = plane_name.clone();
let binding = if let Some(ref target) = target {
format!(" [[texture({})]]", target.planes[i])
} else {
String::new()
};
args.push(EntryPointArgument {
ty_name,
name,
binding,
init: None,
});
}
let params_ty_name = &self.names
[&NameKey::Type(module.special_types.external_texture_params.unwrap())];
let params_name = &self.names[&NameKey::ExternalTextureGlobalVariable(
handle,
ExternalTextureNameKey::Params,
)];
let binding = if let Some(ref target) = target {
format!(" [[buffer({})]]", target.params)
} else {
String::new()
};
// Iterate vbm_resolved, output one argument for every vertex buffer, // using the names we generated earlier. for vbm in &vbm_resolved {
let id = &vbm.id;
let ty_name = &vbm.ty_name;
let param_name = &vbm.param_name;
args.push(EntryPointArgument {
ty_name: format!("const device {ty_name}*"),
name: param_name.clone(),
binding: format!(" [[buffer({id})]]"),
init: None,
});
}
}
// If this entry uses any variable-length arrays, their sizes are // passed as a final struct-typed argument. if needs_buffer_sizes { // this is checked earlier
let resolved = options.resolve_sizes_buffer(ep).unwrap();
let mut binding = String::new();
resolved.try_fmt(&mut binding)?;
args.push(EntryPointArgument {
ty_name: "constant _mslBufferSizes&".to_string(),
name: "_buffer_sizes".to_string(),
binding,
init: None,
});
}
let mut is_first_arg = true; for arg in &args { if is_first_arg {
write!(self.out, " ")?;
} else {
write!(self.out, ", ")?;
}
is_first_arg = false;
write!(self.out, "{} {}", arg.ty_name, arg.name)?; if !is_wrapped {
write!(self.out, "{}", arg.binding)?; if let Some(init) = arg.init {
write!(self.out, " = ")?;
self.put_const_expression(
init,
module,
mod_info,
&module.global_expressions,
)?;
}
}
writeln!(self.out)?;
} if ep.stage == crate::ShaderStage::Mesh { for (handle, var) in module.global_variables.iter() { ifvar.space != crate::AddressSpace::WorkGroup || fun_info[handle].is_empty() { continue;
} if is_first_arg {
write!(self.out, " ")?;
} else {
write!(self.out, ", ")?;
}
let ty_context = TypeContext {
handle: module.global_variables[handle].ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: false,
};
writeln!(
self.out, "threadgroup {ty_context}& {}",
self.names[&NameKey::GlobalVariable(handle)]
)?;
}
}
// end of the entry point argument list
writeln!(self.out, ") {{")?;
// Starting the function body. if do_vertex_pulling { // Provide zero values for all the attributes, which we will overwrite with // real data from the vertex attribute buffers, if the indices are in-bounds. for vbm in &vbm_resolved { for attribute in vbm.attributes {
let location = attribute.shader_location;
let am_option = am_resolved.get(&location); if am_option.is_none() { // This bound attribute isn't used in this entry point, so // don't bother zero-initializing it. continue;
}
let am = am_option.unwrap();
let attribute_ty_name = &am.ty_name;
let attribute_name = &am.name;
// Output a bounds check block that will set real values for the // attributes, if the bounds are satisfied.
write!(self.out, "{}if (", back::Level(1))?;
let idx = &vbm.id;
let stride = &vbm.stride;
let index_name = match vbm.step_mode {
back::msl::VertexBufferStepMode::Constant => "0",
back::msl::VertexBufferStepMode::ByVertex => { if let Some(ref name) = v_existing_id {
name
} else {
&v_id
}
}
back::msl::VertexBufferStepMode::ByInstance => { if let Some(ref name) = i_existing_id {
name
} else {
&i_id
}
}
};
write!(
self.out, "{index_name} < (_buffer_sizes.buffer_size{idx} / {stride})"
)?;
writeln!(self.out, ") {{")?;
// Pull the bytes out of the vertex buffer.
let ty_name = &vbm.ty_name;
let elem_name = &vbm.elem_name;
let param_name = &vbm.param_name;
// Now set real values for each of the attributes, by unpacking the data // from the buffer elements. for attribute in vbm.attributes {
let location = attribute.shader_location;
let Some(am) = am_resolved.get(&location) else { // This bound attribute isn't used in this entry point, so // don't bother extracting the data. Too bad we emitted the // unpacking function earlier -- it might not get used. continue;
};
let attribute_name = &am.name;
let attribute_ty_name = &am.ty_name;
let offset = attribute.offset;
let func = unpacking_functions
.get(&attribute.format)
.expect("Should have generated this unpacking function earlier.");
let func_name = &func.name;
// Check dimensionality of the attribute compared to the unpacking // function. If attribute dimension > unpack dimension, we have to // pad out the unpack value from a vec4(0, 0, 0, 1) of matching // scalar type. Otherwise, if attribute dimension is < unpack // dimension, then we need to explicitly truncate the result.
let needs_padding_or_truncation = am.dimension.cmp(&func.dimension);
// We need an extra type conversion if the shader type does not // match the type returned from the unpacking function.
let needs_conversion = am.scalar != func.scalar;
if needs_padding_or_truncation != Ordering::Equal { // Emit a comment flagging that a conversion is happening, // since the actual logic can be at the end of a long line.
writeln!(
self.out, "{}// {attribute_ty_name} <- {:?}",
back::Level(2),
attribute.format
)?;
}
if needs_padding_or_truncation == Ordering::Greater { // Needs padding: emit constructor call for wider type
write!(self.out, "{attribute_ty_name}(")?;
}
// Emit call to unpacking function if needs_conversion {
put_numeric_type(&mut self.out, am.scalar, func.dimension.as_slice())?;
write!(self.out, "(")?;
}
write!(self.out, "{func_name}({elem_name}.data[{offset}]")?; for i in (offset + 1)..(offset + func.byte_count) {
write!(self.out, ", {elem_name}.data[{i}]")?;
}
write!(self.out, ")")?; if needs_conversion {
write!(self.out, ")")?;
}
match needs_padding_or_truncation {
Ordering::Greater => { // Padding
let ty_is_int = scalar_is_int(am.scalar);
let zero_value = if ty_is_int { "0" } else { "0.0" };
let one_value = if ty_is_int { "1" } else { "1.0" }; for i in func.dimension.map_or(1, u8::from)
..am.dimension.map_or(1, u8::from)
{
write!(
self.out, ", {}", if i == 3 { one_value } else { zero_value }
)?;
}
}
Ordering::Less => { // Truncate to the first `am.dimension` components
write!(
self.out, ".{}",
&"xyzw"[0..usize::from(am.dimension.map_or(1, u8::from))]
)?;
}
Ordering::Equal => {}
}
if needs_padding_or_truncation == Ordering::Greater {
write!(self.out, ")")?;
}
writeln!(self.out, ";")?;
}
// End the bounds check / attribute setting block.
writeln!(self.out, "{}}}", back::Level(1))?;
}
}
// Metal doesn't support private mutable variables outside of functions, // so we put them here, just like the locals. for (handle, var) in module.global_variables.iter() {
let usage = fun_info[handle]; if usage.is_empty() { continue;
} ifvar.space == crate::AddressSpace::Private {
let tyvar = TypedGlobalVariable {
module,
names: &self.names,
handle,
usage,
reference: false,
};
write!(self.out, "{}", back::INDENT)?;
tyvar.try_fmt(&mut self.out)?;
match var.init {
Some(value) => {
write!(self.out, " = ")?;
self.put_const_expression(
value,
module,
mod_info,
&module.global_expressions,
)?;
writeln!(self.out, ";")?;
}
None => {
writeln!(self.out, " = {{}};")?;
}
};
} elseif let Some(ref binding) = var.binding {
let resolved = options.resolve_resource_binding(ep, binding).unwrap(); if let Some(sampler) = resolved.as_inline_sampler(options) { // write an inline sampler
let name = &self.names[&NameKey::GlobalVariable(handle)];
writeln!(
self.out, "{}constexpr {}::sampler {}(",
back::INDENT,
NAMESPACE,
name
)?;
self.put_inline_sampler_properties(back::Level(2), sampler)?;
writeln!(self.out, "{});", back::INDENT)?;
} elseif let crate::TypeInner::Image { class: crate::ImageClass::External,
..
} = module.types[var.ty].inner
{ // Wrap the individual arguments for each external texture global // in a struct which can be easily passed around.
let wrapper_name = &self.names[&NameKey::GlobalVariable(handle)];
let l1 = back::Level(1);
let l2 = l1.next();
writeln!(
self.out, "{l1}const {EXTERNAL_TEXTURE_WRAPPER_STRUCT} {wrapper_name} {{"
)?; for i in 0..3 {
let plane_name = &self.names[&NameKey::ExternalTextureGlobalVariable(
handle,
ExternalTextureNameKey::Plane(i),
)];
writeln!(self.out, "{l2}.plane{i} = {plane_name},")?;
}
let params_name = &self.names[&NameKey::ExternalTextureGlobalVariable(
handle,
ExternalTextureNameKey::Params,
)];
writeln!(self.out, "{l2}.params = {params_name},")?;
writeln!(self.out, "{l1}}};")?;
}
}
}
// Now take the arguments that we gathered into structs, and the // structs that we flattened into arguments, and emit local // variables with initializers that put everything back the way the // body code expects. // // If we had to generate fresh names for struct members passed as // arguments, be sure to use those names when rebuilding the struct. // // "Each day, I change some zeros to ones, and some ones to zeros. // The rest, I leave alone." for (arg_index, arg) in fun.arguments.iter().enumerate() {
let arg_name =
&self.names[&NameKey::EntryPointArgument(ep_index as _, arg_index as u32)];
match module.types[arg.ty].inner {
crate::TypeInner::Struct { ref members, .. } => {
let struct_name = &self.names[&NameKey::Type(arg.ty)];
write!(
self.out, "{}const {} {} = {{ ",
back::INDENT,
struct_name,
arg_name
)?; for (member_index, member) in members.iter().enumerate() {
let key = NameKey::StructMember(arg.ty, member_index as u32);
let name = &flattened_member_names[&key]; if member_index != 0 {
write!(self.out, ", ")?;
} // insert padding initialization, if needed if self
.struct_member_pads
.contains(&(arg.ty, member_index as u32))
{
write!(self.out, "{{}}, ")?;
}
match member.binding {
Some(crate::Binding::Location {
interpolation: Some(crate::Interpolation::PerVertex),
..
}) => {
writeln!(
self.out, "{0}{{ {1}.{2}.get({NAMESPACE}::vertex_index::first), {1}.{2}.get({NAMESPACE}::vertex_index::second), {1}.{2}.get({NAMESPACE}::vertex_index::third) }}",
back::INDENT,
varyings_member_name,
arg_name,
)?; continue;
}
Some(crate::Binding::Location { .. }) => { if has_varyings {
write!(self.out, "{varyings_member_name}.")?;
}
}
_ => (),
}
write!(self.out, "{name}")?;
}
writeln!(self.out, " }};")?;
}
_ => match arg.binding {
Some(crate::Binding::Location {
interpolation: Some(crate::Interpolation::PerVertex),
..
}) => {
let ty_name = TypeContext {
handle: arg.ty,
gctx: module.to_ctx(),
names: &self.names,
access: crate::StorageAccess::empty(),
first_time: false,
};
writeln!(
self.out, "{0}const {ty_name} {arg_name} = {{ {1}.{2}.get({NAMESPACE}::vertex_index::first), {1}.{2}.get({NAMESPACE}::vertex_index::second), {1}.{2}.get({NAMESPACE}::vertex_index::third) }};",
back::INDENT,
varyings_member_name,
arg_name,
)?;
}
Some(crate::Binding::Location { .. })
| Some(crate::Binding::BuiltIn(crate::BuiltIn::Barycentric { .. })) => { if has_varyings {
writeln!(
self.out, "{}const auto {} = {}.{};",
back::INDENT,
arg_name,
varyings_member_name,
arg_name
)?;
}
}
_ => {}
},
}
}
let guarded_indices =
index::find_checked_indexes(module, fun, fun_info, options.bounds_check_policies);
// Finally, declare all the local variables that we need //TODO: we can postpone this till the relevant expressions are emitted
self.put_locals(&context.expression)?;
self.update_expressions_to_bake(fun, fun_info, &context.expression);
self.put_block(back::Level(1), &fun.body, &context)?;
writeln!(self.out, "}}")?; if ep_index + 1 != module.entry_points.len() {
writeln!(self.out)?;
}
self.named_expressions.clear();
pub(super) fn write_barrier(
&mut self,
flags: crate::Barrier,
level: back::Level,
) -> BackendResult { // Note: OR-ring bitflags requires `__HAVE_MEMFLAG_OPERATORS__`, // so we try to avoid it here. if flags.is_empty() {
writeln!(
self.out, "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_none);",
)?;
} if flags.contains(crate::Barrier::STORAGE) {
writeln!(
self.out, "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_device);",
)?;
} if flags.contains(crate::Barrier::WORK_GROUP) {
writeln!(
self.out, "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_threadgroup);",
)?; if self.needs_object_memory_barriers {
writeln!(
self.out, "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_object_data);",
)?;
}
} if flags.contains(crate::Barrier::SUB_GROUP) {
writeln!(
self.out, "{level}{NAMESPACE}::simdgroup_barrier({NAMESPACE}::mem_flags::mem_threadgroup);",
)?;
} if flags.contains(crate::Barrier::TEXTURE) {
writeln!(
self.out, "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_texture);",
)?;
}
Ok(())
}
}
/// Initializing workgroup variables is more tricky for Metal because we have to deal /// with atomics at the type-level (which don't have a copy constructor).
mod workgroup_mem_init {
use crate::EntryPoint;
¤ 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.0.616Bemerkung:
¤
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.
