Add -cuda-relaxed-constexpr, which lets CUDA device code call constexpr functions.

[jlebar]

Summary

In particular, this lets device code call most std::complex functions.

[jlebar]

Actually, this isn't enough. It works fine for std::complex, but it screws with our existing std math business, because some of the host functions that we're trying to override are constexpr, and now they're treated as constexpr host-device.

[jlebar]

Abandoning this for now -- it is a maze of twisty passages all alike. Some std math functions are constexpr, and making them host+device affects all our existing math overloads in fun and exciting ways.

We may need this at some point, but for now I'm going to try something more finely targeted specifically at std::complex.

include/clang/Basic/LangOptions.def

 LANGOPT(OpenMPUseTLS      , 1, 0, "Use TLS for threadprivates or runtime calls")
 LANGOPT(OpenMPIsDevice    , 1, 0, "Generate code only for OpenMP target device")
 
 LANGOPT(CUDAIsDevice      , 1, 0, "Compiling for CUDA device")
 LANGOPT(CUDAAllowHostCallsFromHostDevice, 1, 0, "Allow host device functions to call host functions")
 LANGOPT(CUDADisableTargetCallChecks, 1, 0, "Disable checks for call targets (host, device, etc.)")
 LANGOPT(CUDATargetOverloads, 1, 0, "Enable function overloads based on CUDA target attributes")
 LANGOPT(CUDAAllowVariadicFunctions, 1, 0, "Allow variadic functions in CUDA device code")
+LANGOPT(CUDARelaxedConstexpr, 1, 0, "Allow calls to constexpr functions from CUDA device code")
 
 LANGOPT(AssumeSaneOperatorNew , 1, 1, "implicit __attribute__((malloc)) for C++'s new operators")
 LANGOPT(SizedDeallocation , 1, 0, "enable sized deallocation functions")
 LANGOPT(ConceptsTS , 1, 0, "enable C++ Extensions for Concepts")
 BENIGN_LANGOPT(ElideConstructors , 1, 1, "C++ copy constructor elision")
 BENIGN_LANGOPT(DumpRecordLayouts , 1, 0, "dumping the layout of IRgen'd records")
 BENIGN_LANGOPT(DumpRecordLayoutsSimple , 1, 0, "dumping the layout of IRgen'd records in a simple form")
 BENIGN_LANGOPT(DumpVTableLayouts , 1, 0, "dumping the layouts of emitted vtables")

include/clang/Driver/CC1Options.td

def fcuda_disable_target_call_checks : Flag<["-"],
     "fcuda-disable-target-call-checks">,
   HelpText<"Disable all cross-target (host, device, etc.) call checks in CUDA">;
 def fcuda_include_gpubinary : Separate<["-"], "fcuda-include-gpubinary">,
   HelpText<"Incorporate CUDA device-side binary into host object file.">;
 def fcuda_target_overloads : Flag<["-"], "fcuda-target-overloads">,
   HelpText<"Enable function overloads based on CUDA target attributes.">;
 def fcuda_allow_variadic_functions : Flag<["-"], "fcuda-allow-variadic-functions">,
   HelpText<"Allow variadic functions in CUDA device code.">;
+def fcuda_relaxed_constexpr : Flag<["-"], "fcuda-relaxed-constexpr">,
+  HelpText<"Allow calls to constexpr functions from CUDA device code.">;
 
 //===----------------------------------------------------------------------===//
 // OpenMP Options
 //===----------------------------------------------------------------------===//
 
 def fopenmp_is_device : Flag<["-"], "fopenmp-is-device">,
   HelpText<"Generate code only for an OpenMP target device.">;
 def fomp_host_ir_file_path : Separate<["-"], "fomp-host-ir-file-path">,

include/clang/Driver/Options.td

 def cuda_gpu_arch_EQ : Joined<["--"], "cuda-gpu-arch=">,
   Flags<[DriverOption, HelpHidden]>, HelpText<"CUDA GPU architecture">;
 def cuda_host_only : Flag<["--"], "cuda-host-only">,
   HelpText<"Do host-side CUDA compilation only">;
 def cuda_noopt_device_debug : Flag<["--"], "cuda-noopt-device-debug">,
   HelpText<"Enable device-side debug info generation. Disables ptxas optimizations.">;
 def cuda_path_EQ : Joined<["--"], "cuda-path=">, Group<i_Group>,
   HelpText<"CUDA installation path">;
+def cuda_relaxed_constexpr : Flag<["--"], "cuda-relaxed-constexpr">,
+  HelpText<"Allow calls to constexpr functions from CUDA device code">;
 def dA : Flag<["-"], "dA">, Group<d_Group>;
 def dD : Flag<["-"], "dD">, Group<d_Group>, Flags<[CC1Option]>,
   HelpText<"Print macro definitions in -E mode in addition to normal output">;
 def dM : Flag<["-"], "dM">, Group<d_Group>, Flags<[CC1Option]>,
   HelpText<"Print macro definitions in -E mode instead of normal output">;
 def dead__strip : Flag<["-"], "dead_strip">;
 def dependency_file : Separate<["-"], "dependency-file">, Flags<[CC1Option]>,
   HelpText<"Filename (or -) to write dependency output to">;

lib/Driver/Tools.cpp

else if (&getToolChain() == C.getCudaHostToolChain())
       AuxToolChain = C.getCudaDeviceToolChain();
     else
       llvm_unreachable("Can't figure out CUDA compilation mode.");
     assert(AuxToolChain != nullptr && "No aux toolchain.");
     CmdArgs.push_back("-aux-triple");
     CmdArgs.push_back(Args.MakeArgString(AuxToolChain->getTriple().str()));
     CmdArgs.push_back("-fcuda-target-overloads");
     CmdArgs.push_back("-fcuda-disable-target-call-checks");
+
+    if (Args.hasArg(options::OPT_cuda_relaxed_constexpr))
+      CmdArgs.push_back("-fcuda-relaxed-constexpr");
   }
 
   if (Triple.isOSWindows() && (Triple.getArch() == llvm::Triple::arm ||
                                Triple.getArch() == llvm::Triple::thumb)) {
     unsigned Offset = Triple.getArch() == llvm::Triple::arm ? 4 : 6;
     unsigned Version;
     Triple.getArchName().substr(Offset).getAsInteger(10, Version);
     if (Version < 7)

lib/Frontend/CompilerInvocation.cpp

if (Args.hasArg(OPT_fcuda_disable_target_call_checks))
     Opts.CUDADisableTargetCallChecks = 1;
 
   if (Args.hasArg(OPT_fcuda_target_overloads))
     Opts.CUDATargetOverloads = 1;
 
   if (Args.hasArg(OPT_fcuda_allow_variadic_functions))
     Opts.CUDAAllowVariadicFunctions = 1;
 
+  if (Args.hasArg(OPT_fcuda_relaxed_constexpr))
+    Opts.CUDARelaxedConstexpr = 1;
+
   if (Opts.ObjC1) {
     if (Arg *arg = Args.getLastArg(OPT_fobjc_runtime_EQ)) {
       StringRef value = arg->getValue();
       if (Opts.ObjCRuntime.tryParse(value))
         Diags.Report(diag::err_drv_unknown_objc_runtime) << value;
     }
 
     if (Args.hasArg(OPT_fobjc_gc_only))

lib/Sema/SemaCUDA.cpp

if (D->hasAttr<CUDAHostAttr>())
       return CFT_HostDevice;
     return CFT_Device;
   } else if (D->hasAttr<CUDAHostAttr>()) {
     return CFT_Host;
   } else if (D->isImplicit()) {
     // Some implicit declarations (like intrinsic functions) are not marked.
     // Set the most lenient target on them for maximal flexibility.
     return CFT_HostDevice;
+  } else if (getLangOpts().CUDARelaxedConstexpr && D->isConstexpr()) {
+    return CFT_HostDevice;
   }
 
   return CFT_Host;
 }
 
 // * CUDA Call preference table
 //
 // F - from,

test/SemaCUDA/relaxed-constexpr.cu

+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify -verify-ignore-unexpected=note -DERR %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fcuda-is-device -verify -verify-ignore-unexpected=note -DERR %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fcuda-relaxed-constexpr -verify -verify-ignore-unexpected=note %s
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -fcuda-is-device -fcuda-relaxed-constexpr -verify -verify-ignore-unexpected=note %s
+
+#include "Inputs/cuda.h"
+
+constexpr int foo() { return 0; }
+
+__host__ void host() {
+  constexpr auto x = foo();
+}
+
+__device__ void device() {
+  constexpr auto x = foo();
+#ifdef ERR
+  // expected-error@-2 {{no matching function for call to 'foo'}}
+#else
+  // expected-no-diagnostics
+#endif
+}