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_cg/src/gpu.rs

@@ -1383,15 +1383,15 @@ extern \"C\" {} {}(",
                     let cg_tile = self.get_cg_tile(nesting_fork.unwrap(), CGType::Use);
                     #[allow(unreachable_patterns)]
                     let cg_op = match intrinsic {
-                        Intrinsic::Max => "max",
-                        Intrinsic::Min => "min",
+                        Intrinsic::Max => "greater",
+                        Intrinsic::Min => "less",
                         _ => unreachable!(),
                     };
                     let id_type_name = self.get_type(id_type, false);
                     write!(
                         w,
                         "{}{} = cg::reduce({}, {}, cg::{}<{}>());\n",
-                        tabs, define_variable, non_reduce_arg, cg_tile, cg_op, id_type_name
+                        tabs, define_variable, cg_tile, non_reduce_arg, cg_op, id_type_name
                     )?;
                 } else {
                     let ty = &self.types[id_type.idx()];