diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index e0b3d31e8..a7374915c 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -376,6 +376,11 @@ pub struct Writer { /// 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, u32)>, + + /// Name of the loop reachability macro. + /// + /// See `emit_loop_reachable_macro` for details. + loop_reachable_macro_name: String, } impl crate::Scalar { @@ -665,6 +670,7 @@ impl Writer { #[cfg(test)] put_block_stack_pointers: Default::default(), struct_member_pads: FastHashSet::default(), + loop_reachable_macro_name: String::default(), } } @@ -675,6 +681,125 @@ impl Writer { self.out } + /// Define a macro to invoke before loops, to defeat MSL infinite loop + /// reasoning. + /// + /// If we haven't done so already, emit the definition of a preprocessor + /// macro to be invoked before each loop 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: + /// + /// ```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. + /// + /// Ideally, Naga would prevent UB entirely via some means that persuades + /// 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. Further experimentation hasn't produced a + /// satisfactory approach. + /// + /// Instead, we accept that the MSL compiler may determine that some loops + /// are infinite, and focus instead on preventing the range analysis from + /// being affected. We transform *every* loop into something like this: + /// + /// ```ignore + /// if (volatile bool unpredictable = true; unpredictable) + /// while (true) { } + /// ``` + /// + /// Since the `volatile` qualifier prevents the compiler from assuming that + /// the `if` condition is true, it cannot be sure the infinite loop is + /// reached, and thus it cannot assume the entire structure is unreachable. + /// 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, before each 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. + /// + /// To make our output a bit more legible, we pull the condition out into a + /// preprocessor macro defined at the top of the module. + fn emit_loop_reachable_macro(&mut self) -> BackendResult { + if !self.loop_reachable_macro_name.is_empty() { + return Ok(()); + } + + self.loop_reachable_macro_name = self.namer.call("LOOP_IS_REACHABLE"); + let loop_reachable_volatile_name = self.namer.call("unpredictable_jump_over_loop"); + writeln!( + self.out, + "#define {} if (volatile bool {} = true; {})", + self.loop_reachable_macro_name, + loop_reachable_volatile_name, + loop_reachable_volatile_name, + )?; + + Ok(()) + } + fn put_call_parameters( &mut self, parameters: impl Iterator>, @@ -2924,10 +3049,15 @@ impl Writer { ref continuing, break_if, } => { + self.emit_loop_reachable_macro()?; if !continuing.is_empty() || break_if.is_some() { let gate_name = self.namer.call("loop_init"); writeln!(self.out, "{level}bool {gate_name} = true;")?; - writeln!(self.out, "{level}while(true) {{")?; + writeln!( + self.out, + "{level}{} while(true) {{", + self.loop_reachable_macro_name, + )?; let lif = level.next(); let lcontinuing = lif.next(); writeln!(self.out, "{lif}if (!{gate_name}) {{")?; @@ -2942,7 +3072,11 @@ impl Writer { writeln!(self.out, "{lif}}}")?; writeln!(self.out, "{lif}{gate_name} = false;")?; } else { - writeln!(self.out, "{level}while(true) {{")?; + writeln!( + self.out, + "{level}{} while(true) {{", + self.loop_reachable_macro_name, + )?; } self.put_block(level.next(), body, context)?; writeln!(self.out, "{level}}}")?; @@ -3379,6 +3513,7 @@ impl Writer { &[CLAMPED_LOD_LOAD_PREFIX], &mut self.names, ); + self.loop_reachable_macro_name.clear(); self.struct_member_pads.clear(); writeln!( diff --git a/naga/tests/out/msl/boids.msl b/naga/tests/out/msl/boids.msl index ce1ccc7cc..0dd520ac7 100644 --- a/naga/tests/out/msl/boids.msl +++ b/naga/tests/out/msl/boids.msl @@ -55,8 +55,9 @@ kernel void main_( vPos = _e8; metal::float2 _e14 = particlesSrc.particles[index].vel; vVel = _e14; +#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop) bool loop_init = true; - while(true) { + LOOP_IS_REACHABLE while(true) { if (!loop_init) { uint _e91 = i; i = _e91 + 1u; diff --git a/naga/tests/out/msl/break-if.msl b/naga/tests/out/msl/break-if.msl index 8c0d9343b..3684f7222 100644 --- a/naga/tests/out/msl/break-if.msl +++ b/naga/tests/out/msl/break-if.msl @@ -7,8 +7,9 @@ using metal::uint; void breakIfEmpty( ) { +#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop) bool loop_init = true; - while(true) { + LOOP_IS_REACHABLE while(true) { if (!loop_init) { if (true) { break; @@ -25,7 +26,7 @@ void breakIfEmptyBody( bool b = {}; bool c = {}; bool loop_init_1 = true; - while(true) { + LOOP_IS_REACHABLE while(true) { if (!loop_init_1) { b = a; bool _e2 = b; @@ -46,7 +47,7 @@ void breakIf( bool d = {}; bool e = {}; bool loop_init_2 = true; - while(true) { + LOOP_IS_REACHABLE while(true) { if (!loop_init_2) { bool _e5 = e; if (a_1 == e) { @@ -65,7 +66,7 @@ void breakIfSeparateVariable( ) { uint counter = 0u; bool loop_init_3 = true; - while(true) { + LOOP_IS_REACHABLE while(true) { if (!loop_init_3) { uint _e5 = counter; if (counter == 5u) { diff --git a/naga/tests/out/msl/collatz.msl b/naga/tests/out/msl/collatz.msl index e28374145..1ae910de6 100644 --- a/naga/tests/out/msl/collatz.msl +++ b/naga/tests/out/msl/collatz.msl @@ -19,7 +19,8 @@ uint collatz_iterations( uint n = {}; uint i = 0u; n = n_base; - while(true) { +#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop) + LOOP_IS_REACHABLE while(true) { uint _e4 = n; if (_e4 > 1u) { } else { diff --git a/naga/tests/out/msl/control-flow.msl b/naga/tests/out/msl/control-flow.msl index 11771693a..dbf75163a 100644 --- a/naga/tests/out/msl/control-flow.msl +++ b/naga/tests/out/msl/control-flow.msl @@ -31,7 +31,8 @@ void switch_case_break( void loop_switch_continue( int x ) { - while(true) { +#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop) + LOOP_IS_REACHABLE while(true) { switch(x) { case 1: { continue; @@ -49,7 +50,7 @@ void loop_switch_continue_nesting( int y, int z ) { - while(true) { + LOOP_IS_REACHABLE while(true) { switch(x_1) { case 1: { continue; @@ -60,7 +61,7 @@ void loop_switch_continue_nesting( continue; } default: { - while(true) { + LOOP_IS_REACHABLE while(true) { switch(z) { case 1: { continue; @@ -85,7 +86,7 @@ void loop_switch_continue_nesting( } } } - while(true) { + LOOP_IS_REACHABLE while(true) { switch(y) { case 1: default: { @@ -108,7 +109,7 @@ void loop_switch_omit_continue_variable_checks( int w ) { int pos_1 = 0; - while(true) { + LOOP_IS_REACHABLE while(true) { switch(x_2) { case 1: { pos_1 = 1; @@ -119,7 +120,7 @@ void loop_switch_omit_continue_variable_checks( } } } - while(true) { + LOOP_IS_REACHABLE while(true) { switch(x_2) { case 1: { break; diff --git a/naga/tests/out/msl/do-while.msl b/naga/tests/out/msl/do-while.msl index c1b4d08b0..b093da1dc 100644 --- a/naga/tests/out/msl/do-while.msl +++ b/naga/tests/out/msl/do-while.msl @@ -8,8 +8,9 @@ using metal::uint; void fb1_( thread bool& cond ) { +#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop) bool loop_init = true; - while(true) { + LOOP_IS_REACHABLE while(true) { if (!loop_init) { bool _e1 = cond; if (!(cond)) { diff --git a/naga/tests/out/msl/overrides-ray-query.msl b/naga/tests/out/msl/overrides-ray-query.msl index 3a508b6f6..f2ad45c98 100644 --- a/naga/tests/out/msl/overrides-ray-query.msl +++ b/naga/tests/out/msl/overrides-ray-query.msl @@ -33,7 +33,8 @@ kernel void main_( rq.intersector.force_opacity((desc.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (desc.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none); rq.intersector.accept_any_intersection((desc.flags & 4) != 0); rq.intersection = rq.intersector.intersect(metal::raytracing::ray(desc.origin, desc.dir, desc.tmin, desc.tmax), acc_struct, desc.cull_mask); rq.ready = true; - while(true) { +#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop) + LOOP_IS_REACHABLE while(true) { bool _e31 = rq.ready; rq.ready = false; if (_e31) { diff --git a/naga/tests/out/msl/ray-query.msl b/naga/tests/out/msl/ray-query.msl index fbdaef548..129ad108a 100644 --- a/naga/tests/out/msl/ray-query.msl +++ b/naga/tests/out/msl/ray-query.msl @@ -53,7 +53,8 @@ RayIntersection query_loop( rq.intersector.force_opacity((_e8.flags & 1) != 0 ? metal::raytracing::forced_opacity::opaque : (_e8.flags & 2) != 0 ? metal::raytracing::forced_opacity::non_opaque : metal::raytracing::forced_opacity::none); rq.intersector.accept_any_intersection((_e8.flags & 4) != 0); rq.intersection = rq.intersector.intersect(metal::raytracing::ray(_e8.origin, _e8.dir, _e8.tmin, _e8.tmax), acs, _e8.cull_mask); rq.ready = true; - while(true) { +#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop) + LOOP_IS_REACHABLE while(true) { bool _e9 = rq.ready; rq.ready = false; if (_e9) { diff --git a/naga/tests/out/msl/shadow.msl b/naga/tests/out/msl/shadow.msl index 3aefc21ee..f8aeef9d4 100644 --- a/naga/tests/out/msl/shadow.msl +++ b/naga/tests/out/msl/shadow.msl @@ -100,8 +100,9 @@ fragment fs_mainOutput fs_main( metal::float3 color = c_ambient; uint i = 0u; metal::float3 normal_1 = metal::normalize(in.world_normal); +#define LOOP_IS_REACHABLE if (volatile bool unpredictable_jump_over_loop = true; unpredictable_jump_over_loop) bool loop_init = true; - while(true) { + LOOP_IS_REACHABLE while(true) { if (!loop_init) { uint _e40 = i; i = _e40 + 1u; @@ -151,7 +152,7 @@ fragment fs_main_without_storageOutput fs_main_without_storage( uint i_1 = 0u; metal::float3 normal_2 = metal::normalize(in_1.world_normal); bool loop_init_1 = true; - while(true) { + LOOP_IS_REACHABLE while(true) { if (!loop_init_1) { uint _e40 = i_1; i_1 = _e40 + 1u;