From 3fda684eb9e69c78b16312a3e927e3ea82e853d1 Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Mon, 16 Sep 2024 12:56:58 -0700 Subject: [PATCH] [naga msl-out] Defeat the MSL compiler's infinite loop analysis. See the comments in the code for details. This patch emits the definition of the macro only when the first loop is encountered. This does make that first loop's code look a bit odd: it would be more natural to define the macro at the top of the file. (See the modified files in `naga/tests/out/msl`.) Rejected alternatives: - We could emit the macro definition unconditionally at the top of the file. But this changes every MSL snapshot output file, whereas only eight of them actually contain loops. - We could have the validator flag modules that contain loops. But the changes end up being not small, and spread across the validator, so this seems disproportionate. If we had other consumers of this information, it might make sense. - We could change the MSL backend to allow text to be generated out of order, so that we can decide whether to define the macro after we've generated all the function bodies. But at the moment this seems like unnecessary complexity, although it might be worth doing in the future if we had additional uses for it - say, to conditionally emit helper function definitions. Fixes #4972. --- naga/src/back/msl/writer.rs | 139 ++++++++++++++++++++- naga/tests/out/msl/boids.msl | 3 +- naga/tests/out/msl/break-if.msl | 9 +- naga/tests/out/msl/collatz.msl | 3 +- naga/tests/out/msl/control-flow.msl | 13 +- naga/tests/out/msl/do-while.msl | 3 +- naga/tests/out/msl/overrides-ray-query.msl | 3 +- naga/tests/out/msl/ray-query.msl | 3 +- naga/tests/out/msl/shadow.msl | 5 +- 9 files changed, 162 insertions(+), 19 deletions(-) diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index e0b3d31e84..a7374915c9 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 ce1ccc7cc2..0dd520ac74 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 8c0d9343b9..3684f7222c 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 e283741459..1ae910de6f 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 11771693aa..dbf75163aa 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 c1b4d08b0e..b093da1dc5 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 3a508b6f61..f2ad45c985 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 fbdaef5484..129ad108a9 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 3aefc21eee..f8aeef9d45 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;