GPU backend

Adds GPU backend and "CUDA" feature to all the tests. As of MR creation, there were no forks in tested IRs so only vanilla non-parallel (single block, single thread) codegen has been tested.

hercules_cg/src/gpu.rs

@@ -470,6 +470,12 @@ namespace cg = cooperative_groups;
                 DynamicConstant::Rem(left, right) => {
                     write!(w, "\t{} = dc{} % dc{};\n", dc_val, left.idx(), right.idx())?
                 }
+                DynamicConstant::Min(left, right) => {
+                    write!(w, "\t{} = min(dc{}, dc{});\n", dc_val, left.idx(), right.idx())?
+                }
+                DynamicConstant::Max(left, right) => {
+                    write!(w, "\t{} = max(dc{}, dc{});\n", dc_val, left.idx(), right.idx())?
+                }
             }
         }
         Ok(())
@@ -1880,8 +1886,8 @@ extern \"C\" {} {}(", if ret_primitive { ret_type.clone() } else { "void".to_str
                 _ => "max",
             },
             Intrinsic::Min => match ty {
-                Type::Float32 => "__fminf",
-                Type::Float64 => "__fmin",
+                Type::Float32 => "fminf",
+                Type::Float64 => "fmin",
                 ty if ty.is_signed() => "smin",
                 ty if ty.is_unsigned() => "umin",
                 _ => "min",