Struct Writer

pub struct Writer<W> {
    out: W,
    names: FastHashMap<NameKey, String>,
    named_expressions: FastIndexMap<Handle<Expression>, String>,
    need_bake_expressions: NeedBakeExpressions,
    namer: Namer,
    struct_member_pads: FastHashSet<(Handle<Type>, u32)>,
    force_bounded_loop_macro_name: String,
}

Fields

out: W

names: FastHashMap<NameKey, String>

named_expressions: FastIndexMap<Handle<Expression>, String>

need_bake_expressions: NeedBakeExpressions

Set of expressions that need to be baked to avoid unnecessary repetition in output

namer: Namer

struct_member_pads: FastHashSet<(Handle<Type>, u32)>

Set of (struct type, struct field index) denoting which fields require padding inserted before them (i.e. between fields at index - 1 and index)

force_bounded_loop_macro_name: String

Name of the force-bounded-loop macro.

See emit_force_bounded_loop_macro for details.

Implementations

impl<W: Write> Writer<W>

pub(super) fn need_workgroup_variables_initialization( &mut self, options: &Options, ep: &EntryPoint, module: &Module, fun_info: &FunctionInfo, ) -> bool

pub(super) fn write_workgroup_variables_initialization( &mut self, module: &Module, module_info: &ModuleInfo, fun_info: &FunctionInfo, local_invocation_id: Option<&NameKey>, ) -> Result<(), Error>

fn write_workgroup_variable_initialization( &mut self, module: &Module, module_info: &ModuleInfo, ty: Handle<Type>, access_stack: &mut AccessStack, level: Level, ) -> Result<(), Error>

impl<W: Write> Writer<W>

pub fn new(out: W) -> Self

Creates a new Writer instance.

pub fn finish(self) -> W

Finishes writing and returns the output.

fn emit_force_bounded_loop_macro(&mut self) -> Result<(), Error>

Define a macro to invoke at the bottom of each loop body, to defeat MSL infinite loop reasoning.

If we haven’t done so already, emit the definition of a preprocessor macro to be invoked at the end of each loop body in the generated MSL, to ensure that the MSL compiler’s optimizations do not remove bounds checks.

Only the first call to this function for a given module actually causes the macro definition to be written. Subsequent loops can simply use the prior macro definition, since macros aren’t block-scoped.

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:

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:

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:

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.

Instead, we add the following code to the bottom of every loop:

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, what makes this a kludge, not a hack, is that this solution leaves the GPU executing a pointless conditional branch, at runtime, in every iteration of the loop. There’s no part of the system that has a global enough view to be sure that unpredictable is true, and remove it from the code. Adding the branch also affects optimization: for example, it’s impossible to unroll this loop. This transformation has been observed to significantly hurt performance.

To make our output a bit more legible, we pull the condition out into a preprocessor macro defined at the top of the module.

This approach is also used by Chromium WebGPU’s Dawn shader compiler: https://dawn.googlesource.com/dawn/+/a37557db581c2b60fb1cd2c01abdb232927dd961/src/tint/lang/msl/writer/printer/printer.cc#222

fn put_call_parameters( &mut self, parameters: impl Iterator<Item = Handle<Expression>>, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_call_parameters_impl<C, E>( &mut self, parameters: impl Iterator<Item = Handle<Expression>>, ctx: &C, put_expression: E, ) -> Result<(), Error>

where E: Fn(&mut Self, &C, Handle<Expression>) -> Result<(), Error>,

fn put_level_of_detail( &mut self, level: LevelOfDetail, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_image_query( &mut self, image: Handle<Expression>, query: &str, level: Option<LevelOfDetail>, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_image_size_query( &mut self, image: Handle<Expression>, level: Option<LevelOfDetail>, kind: ScalarKind, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_cast_to_uint_scalar_or_vector( &mut self, expr: Handle<Expression>, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_image_sample_level( &mut self, image: Handle<Expression>, level: SampleLevel, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_image_coordinate_limits( &mut self, image: Handle<Expression>, level: Option<LevelOfDetail>, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_restricted_scalar_image_index( &mut self, image: Handle<Expression>, index: Handle<Expression>, limit_method: &str, context: &ExpressionContext<'_>, ) -> Result<(), Error>

General function for writing restricted image indexes.

This is used to produce restricted mip levels, array indices, and sample indices for ImageLoad and ImageStore accesses under the Restrict bounds check policy.

This function writes an expression of the form:

    metal::min(uint(INDEX), IMAGE.LIMIT_METHOD() - 1)

fn put_restricted_texel_address( &mut self, image: Handle<Expression>, address: &TexelAddress, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_image_access_bounds_check( &mut self, image: Handle<Expression>, address: &TexelAddress, context: &ExpressionContext<'_>, ) -> Result<(), Error>

Write an expression that is true if the given image access is in bounds.

fn put_image_load( &mut self, load: Handle<Expression>, image: Handle<Expression>, address: TexelAddress, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_unchecked_image_load( &mut self, image: Handle<Expression>, address: &TexelAddress, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_image_atomic( &mut self, level: Level, image: Handle<Expression>, address: &TexelAddress, fun: AtomicFunction, value: Handle<Expression>, context: &StatementContext<'_>, ) -> Result<(), Error>

fn put_image_store( &mut self, level: Level, image: Handle<Expression>, address: &TexelAddress, value: Handle<Expression>, context: &StatementContext<'_>, ) -> Result<(), Error>

fn put_dynamic_array_max_index( &mut self, handle: Handle<GlobalVariable>, context: &ExpressionContext<'_>, ) -> Result<(), Error>

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_dot_product( &mut self, arg: Handle<Expression>, arg1: Handle<Expression>, size: usize, context: &ExpressionContext<'_>, ) -> Result<(), Error>

Emit code for the arithmetic expression of the dot product.

fn put_isign( &mut self, arg: Handle<Expression>, context: &ExpressionContext<'_>, ) -> Result<(), Error>

Emit code for the isign expression.

fn put_const_expression( &mut self, expr_handle: Handle<Expression>, module: &Module, mod_info: &ModuleInfo, ) -> Result<(), Error>

fn put_possibly_const_expression<C, I, E>( &mut self, expr_handle: Handle<Expression>, expressions: &Arena<Expression>, module: &Module, mod_info: &ModuleInfo, ctx: &C, get_expr_ty: I, put_expression: E, ) -> Result<(), Error>

where I: Fn(&C, Handle<Expression>) -> &TypeResolution, E: Fn(&mut Self, &C, Handle<Expression>) -> Result<(), Error>,

fn put_expression( &mut self, expr_handle: Handle<Expression>, context: &ExpressionContext<'_>, is_scoped: bool, ) -> Result<(), Error>

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.

fn put_binop<F>( &mut self, op: BinaryOperator, left: Handle<Expression>, right: Handle<Expression>, context: &ExpressionContext<'_>, is_scoped: bool, put_expression: &F, ) -> Result<(), Error>

where F: Fn(&mut Self, Handle<Expression>, &ExpressionContext<'_>, bool) -> Result<(), Error>,

Emits code for a binary operation, using the provided callback to emit the left and right operands.

fn put_wrapped_expression_for_packed_vec3_access<F>( &mut self, expr_handle: Handle<Expression>, context: &ExpressionContext<'_>, is_scoped: bool, put_expression: &F, ) -> Result<(), Error>

where F: Fn(&mut Self, Handle<Expression>, &ExpressionContext<'_>, bool) -> Result<(), Error>,

Used by expressions like Swizzle and Binary since they need packed_vec3’s to be casted to a vec3

fn put_bitcasted_expression<F>( &mut self, cast_to: &TypeInner, context: &ExpressionContext<'_>, put_expression: &F, ) -> Result<(), Error>

where F: Fn(&mut Self, &ExpressionContext<'_>, bool) -> Result<(), Error>,

Emits code for an expression using the provided callback, wrapping the result in a bitcast to the type cast_to.

fn put_index( &mut self, index: GuardedIndex, context: &ExpressionContext<'_>, is_scoped: bool, ) -> Result<(), Error>

Write a GuardedIndex as a Metal expression.

fn put_bounds_checks( &mut self, chain: Handle<Expression>, context: &ExpressionContext<'_>, level: Level, prefix: &'static str, ) -> Result<bool, Error>

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:

{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.

fn put_access_chain( &mut self, chain: Handle<Expression>, policy: BoundsCheckPolicy, context: &ExpressionContext<'_>, ) -> Result<(), Error>

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.

fn put_subscripted_access_chain( &mut self, base: Handle<Expression>, base_ty: &TypeInner, index: GuardedIndex, policy: BoundsCheckPolicy, context: &ExpressionContext<'_>, ) -> Result<(), Error>

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.

fn put_load( &mut self, pointer: Handle<Expression>, context: &ExpressionContext<'_>, is_scoped: bool, ) -> Result<(), Error>

fn put_unchecked_load( &mut self, pointer: Handle<Expression>, policy: BoundsCheckPolicy, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn put_return_value( &mut self, level: Level, expr_handle: Handle<Expression>, result_struct: Option<&str>, context: &ExpressionContext<'_>, ) -> Result<(), Error>

fn update_expressions_to_bake( &mut self, func: &Function, info: &FunctionInfo, context: &ExpressionContext<'_>, )

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 start_baking_expression( &mut self, handle: Handle<Expression>, context: &ExpressionContext<'_>, name: &str, ) -> Result<(), Error>

fn put_cache_restricted_level( &mut self, load: Handle<Expression>, image: Handle<Expression>, mip_level: Option<Handle<Expression>>, indent: Level, context: &StatementContext<'_>, ) -> Result<(), Error>

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.

fn put_block( &mut self, level: Level, statements: &[Statement], context: &StatementContext<'_>, ) -> Result<(), Error>

fn put_store( &mut self, pointer: Handle<Expression>, value: Handle<Expression>, level: Level, context: &StatementContext<'_>, ) -> Result<(), Error>

fn put_unchecked_store( &mut self, pointer: Handle<Expression>, value: Handle<Expression>, policy: BoundsCheckPolicy, level: Level, context: &StatementContext<'_>, ) -> Result<(), Error>

pub fn write( &mut self, module: &Module, info: &ModuleInfo, options: &Options, pipeline_options: &PipelineOptions, ) -> Result<TranslationInfo, Error>

fn put_default_constructible(&mut self) -> Result<(), Error>

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.

fn put_ray_query_type(&mut self) -> Result<(), Error>

fn write_type_defs(&mut self, module: &Module) -> Result<(), Error>

fn write_global_constants( &mut self, module: &Module, mod_info: &ModuleInfo, ) -> Result<(), Error>

Writes all named constants

fn put_inline_sampler_properties( &mut self, level: Level, sampler: &InlineSampler, ) -> Result<(), Error>

fn write_unpacking_function( &mut self, format: VertexFormat, ) -> Result<(String, u32, u32), Error>

fn write_functions( &mut self, module: &Module, mod_info: &ModuleInfo, options: &Options, pipeline_options: &PipelineOptions, ) -> Result<TranslationInfo, Error>

fn write_barrier(&mut self, flags: Barrier, level: Level) -> Result<(), Error>