Misc. GPU improvements

• GPU backend fixes.
  • Emit float infinity properly.
  • Emit thread_block_tile variables at start of function.
• RT backend fixes.
  • Lower intrinsics.
  • Emit float infinity properly.
• Add max / min to monoid reduction cleaning (utilities for working w/ smallest/largest values of a datatype).
• Don't outline scalar constants (TODO: inline these interprocedurally later).
• Optimize edge detection schedules, especially GPU.
  • Emit a 2-level max reduction tree.

hercules_opt/src/editor.rs

@@ -795,6 +795,26 @@ impl<'a, 'b> FunctionEdit<'a, 'b> {
         self.add_constant(constant_to_construct)
     }
 
+    pub fn add_pos_inf_constant(&mut self, id: TypeID) -> ConstantID {
+        let ty = self.get_type(id).clone();
+        let constant_to_construct = match ty {
+            Type::Float32 => Constant::Float32(ordered_float::OrderedFloat(f32::INFINITY)),
+            Type::Float64 => Constant::Float64(ordered_float::OrderedFloat(f64::INFINITY)),
+            _ => panic!(),
+        };
+        self.add_constant(constant_to_construct)
+    }
+
+    pub fn add_neg_inf_constant(&mut self, id: TypeID) -> ConstantID {
+        let ty = self.get_type(id).clone();
+        let constant_to_construct = match ty {
+            Type::Float32 => Constant::Float32(ordered_float::OrderedFloat(f32::NEG_INFINITY)),
+            Type::Float64 => Constant::Float64(ordered_float::OrderedFloat(f64::NEG_INFINITY)),
+            _ => panic!(),
+        };
+        self.add_constant(constant_to_construct)
+    }
+
     pub fn get_constant(&self, id: ConstantID) -> impl Deref<Target = Constant> + '_ {
         if id.idx() < self.editor.constants.borrow().len() {
             Either::Left(Ref::map(self.editor.constants.borrow(), |constants| {