[CUDA] Make unattributed constexpr functions (usually) implicitly host+device.

[jlebar]

Summary

[CUDA] Make unattributed constexpr functions implicitly host+device.

With this patch, by a constexpr function is implicitly host+device
unless:

a) it's a variadic function (variadic functions are not allowed on the device side), or
b) it's preceeded by a device overload in a system header.

The restriction on overloading host device functions on the
basis of their CUDA attributes remains in place, but we use (b) to allow
us to define device overloads for constexpr functions in cmath,
which would otherwise be host device and thus not overloadable.

You can disable this behavior with -fno-cuda-host-device-constexpr.

[jlebar]

Actually run the tests, and fix the CUDA overloading test.

[tra]

Now that H/D and HD cal all be in the same overload set, we'll also need additional tests in CodeGenCUDA/function-overload.cu for cases that now became legal.

[jlebar]

In D18380#381025, @tra wrote:

Now that H/D and HD cal all be in the same overload set, we'll also need additional tests in CodeGenCUDA/function-overload.cu for cases that now became legal.

There are lots of tests that used to be compile errors and now aren't -- what do you think we're missing?

[tra]

We need tests to demonstrate that we pick correct function when we have mix
of HD+H/D in the overload set.
Existing tests only cover resolution of {HD,HD}, {H,H} {D,D} {H,D} sets

[jlebar]

Add tests checking host+device overloading.

[jlebar]

In D18380#381031, @tra wrote:

We need tests to demonstrate that we pick correct function when we have mix
of HD+H/D in the overload set.
Existing tests only cover resolution of {HD,HD}, {H,H} {D,D} {H,D} sets

Aha, got it. I think adding this is simple given the existing framework -- lmk what you think.

[jlebar]

Update test as discussed -- now we check that we're invoking the correct overloads.

[rsmith]

include/clang/Driver/CC1Options.td

702–703

Is there a better name we can use for this? I don't think this is "relaxed" in any obvious sense. -fcuda-host-device-constexpr or -fcuda-constexpr-on-device might be clearer?

lib/Driver/Tools.cpp

3597

For flags that are enabled by default, we usually have the -cc1 flag be a -fno-* flag. This allows people to use (for instance) clang blah.cu -Xclang -fno-cuda-relaxed-constexpr if necessary.

lib/Sema/SemaOverload.cpp

1132

No parens around == comparisons.

[jlebar]

include/clang/Driver/CC1Options.td

702–703

"relaxed constexpr" is nvidia's term -- do you think it might be helpful to use the same terminology? I understand there's some prior art here, with respect to clang accepting gcc's flags, although the situation here is of course different.

lib/Driver/Tools.cpp

3597

Yeah, Artem and I had a discussion about this yesterday. As you can see, there are two other flags above which are turned on by default -- these also lack -fno variants.

I think it would be good to be consistent here. I'm tempted to add another patch below this one which makes the other two -fno, then we can make this one -fno as well. It seems that convention is to just get rid of the existing non-fno flags, rather than leave both positive and negative versions.

Does that sound OK to you?

[rsmith]

include/clang/Driver/CC1Options.td

702–703

I think it's problematic to use that terminology, as "relaxed constexpr" is also used to describe the C++14 constexpr rules (see n3652).

lib/Driver/Tools.cpp

3597

Yes, that sounds fine.

[jlebar]

include/clang/Driver/CC1Options.td

702–703

Heh, I can't argue with that.

lib/Driver/Tools.cpp

3597

Okay, thank you. After talking to Artem, we're just going to remove those two flags entirely. So after we convert relaxed-constexpr to an fno flag, there should be no changes to this file in this patch.

[jlebar]

Switch to -fno-cuda-host-device-constexpr. Only implicitly add the attributes
on functions which themselves lack host/device attributes. Add more tests.

[jlebar]

Changed as discussed. Please have another look. Thank you for your continued patience here.

[tra]

lib/Sema/SemaDecl.cpp

8012–8014

Can we have constexpr __global__ ?

[jlebar]

Add check for global constexpr functions.

lib/Sema/SemaDecl.cpp

8012–8014

Yikes. We're saved (unless Richard has a tricky counterexample) because kernels must be void and constexpr must not be void. But I'll add a check here anyway.

[jlebar]

Richard, are you happy here?

[rsmith]

The change to allow __host__ __device__ functions to be overloaded with other combinations of target attributes appears to be separable from the constexpr change; please split it out and commit it first.

include/clang/Basic/LangOptions.def

178

This should be a noun phrase -- this string appears in contexts like "support for %0 is enabled" -- so this should be "treating unattributed [...]".

lib/Sema/SemaDecl.cpp

8012–8014

constexpr functions can return void in a couple of different ways (in C++11, if they're template specializations with dependent return types that instantiate to void, and in C++14 there is no restriction on constexpr functions returning void).

[jlebar]

include/clang/Basic/LangOptions.def

178

Thanks. This is fixed in my patch queue, and I will push a change for the other ones as part of this patch queue.

[jlebar]

Okay, just one more patch, D18458, then I think we're good here. (This is split up into two patches in my queue.)

Thanks for your help, Richard.

[jlebar]

Update per changes to patch description. Now a constexpr becomes implicitly HD
unless there's a preceeding device overload.

[jlebar]

Updated as discussed -- please have a look.

[tra]

I wonder if we can find a way to decide whether particular constexpr function should be treated as HD or not without relying on particular order the functions are seen by compiler (or whether they come from system headers).

Right now we're relying on checking overloads of constexpr's function decl once and applying HD attributes based on state of overload set at the point in TU. We then use those attributes during overload resolution.

What if instead of permanently sticking HD attributes on the constexpr function, we instead postpone decision to the point of overload resolution and figure out effective attributes or call preference based on contents of the whole overload set regardless of the order the decls were added to the set.

test/SemaCUDA/host-device-constexpr.cu

30–31 ↗

(On Diff #51868)

"should prevent this"

[jlebar]

In D18380#385240, @tra wrote:

What if instead of permanently sticking HD attributes on the constexpr function, we instead postpone decision to the point of overload resolution and figure out effective attributes or call preference based on contents of the whole overload set regardless of the order the decls were added to the set.

The problem we were trying to prevent by requiring that the __device__ overload come first is:

constexpr int foo();
__device__ void bar() { foo(); }
__device__ int foo();
__device__ void baz() { foo(); }

In this example, we're forced to instantiate both versions of foo() on the device. Being lazy about making the first foo HD doesn't help, because at the time we see bar, it's the only option available.

(Instantiating both foos is a problem if they have the same mangling. And we want them to have the same mangling so we maintain ABI compatibility with nvcc.)

[tra]

LGTM.

[jlebar]

(Just to be clear, I'm waiting on Richard's review here, even though he lg'ed an version of this patch.)

[rsmith]

LGTM

[jlebar]

Thank you all your time here, Art, Reid, and Richard. Fingers crossed we don't have to worry about this again for a while...

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, "Treat constexpr functions as __host__ __device__")
 
 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<"Treat constexpr functions as __host__ __device__.">;
 
 //===----------------------------------------------------------------------===//
 // 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">,

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");
+    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/SemaDecl.cpp

if (UnifySection(CodeSegStack.CurrentValue->getString(),
                          ASTContext::PSF_Read,
                      NewFD))
       NewFD->dropAttr<SectionAttr>();
   }
 
   // Handle attributes.
   ProcessDeclAttributes(S, NewFD, D);
 
+  // With -fcuda-relaxed-constexpr, constexpr functions are treated as
+  // implicitly __host__ __device__.  Device-side variadic functions are not
+  // allowed, so we just treat those as host-only.
+  if (getLangOpts().CUDA && NewFD->isConstexpr() && !NewFD->isVariadic() &&
+      getLangOpts().CUDARelaxedConstexpr) {
+    NewFD->addAttr(CUDAHostAttr::CreateImplicit(Context));
+    NewFD->addAttr(CUDADeviceAttr::CreateImplicit(Context));
+  }
+
   if (getLangOpts().OpenCL) {
     // OpenCL v1.1 s6.5: Using an address space qualifier in a function return
     // type declaration will generate a compilation error.
     unsigned AddressSpace = NewFD->getReturnType().getAddressSpace();
     if (AddressSpace == LangAS::opencl_local ||
         AddressSpace == LangAS::opencl_global ||
         AddressSpace == LangAS::opencl_constant) {
       Diag(NewFD->getLocation(),

lib/Sema/SemaOverload.cpp

bool Sema::IsOverload(FunctionDecl *New, FunctionDecl *Old,
   if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads) {
     CUDAFunctionTarget NewTarget = IdentifyCUDATarget(New),
                        OldTarget = IdentifyCUDATarget(Old);
     if (NewTarget == CFT_InvalidTarget || NewTarget == CFT_Global)
       return false;
 
     assert((OldTarget != CFT_InvalidTarget) && "Unexpected invalid target.");
 
-    // Don't allow mixing of HD with other kinds. This guarantees that
-    // we have only one viable function with this signature on any
-    // side of CUDA compilation .
-    // __global__ functions can't be overloaded based on attribute
-    // difference because, like HD, they also exist on both sides.
-    if ((NewTarget == CFT_HostDevice) || (OldTarget == CFT_HostDevice) ||
-        (NewTarget == CFT_Global) || (OldTarget == CFT_Global))
+    // Don't allow __global__ functions to be overloaded with other functions,
+    // based solely on their CUDA attributes. This guarantees that we have only
+    // one viable function with this signature on any side of CUDA compilation.
+    if ((NewTarget == CFT_Global) || (OldTarget == CFT_Global))
       return false;
 
     // Allow overloading of functions with same signature, but
     // different CUDA target attributes.
     return NewTarget != OldTarget;
   }
 
   // The signatures match; this is not an overload.

test/SemaCUDA/function-overload.cu

 // Host and unattributed functions can't be overloaded.
 __host__ void hh() {} // expected-note {{previous definition is here}}
 void hh() {} // expected-error {{redefinition of 'hh'}}
 
 // H/D overloading is OK.
 __host__ HostReturnTy dh() { return HostReturnTy(); }
 __device__ DeviceReturnTy dh() { return DeviceReturnTy(); }
 
-// H/HD and D/HD are not allowed.
-__host__ __device__ int hdh() { return 0; } // expected-note {{previous definition is here}}
-__host__ int hdh() { return 0; }            // expected-error {{redefinition of 'hdh'}}
-
-__host__ int hhd() { return 0; }            // expected-note {{previous definition is here}}
-__host__ __device__ int hhd() { return 0; } // expected-error {{redefinition of 'hhd'}}
-// expected-warning@-1 {{attribute declaration must precede definition}}
-// expected-note@-3 {{previous definition is here}}
-
-__host__ __device__ int hdd() { return 0; } // expected-note {{previous definition is here}}
-__device__ int hdd() { return 0; }          // expected-error {{redefinition of 'hdd'}}
-
-__device__ int dhd() { return 0; }          // expected-note {{previous definition is here}}
-__host__ __device__ int dhd() { return 0; } // expected-error {{redefinition of 'dhd'}}
-// expected-warning@-1 {{attribute declaration must precede definition}}
-// expected-note@-3 {{previous definition is here}}
+// H/HD and D/HD are also OK.
+__host__ __device__ HostDeviceReturnTy hdh() { return HostDeviceReturnTy(); }
+__host__ HostReturnTy hdh() { return HostReturnTy(); }
+
+__host__ HostReturnTy hhd() { return HostReturnTy(); }
+__host__ __device__ HostDeviceReturnTy hhd() { return HostDeviceReturnTy(); }
+
+__host__ __device__ HostDeviceReturnTy hdd() { return HostDeviceReturnTy(); }
+__device__ DeviceReturnTy hdd() { return DeviceReturnTy(); }
+
+__device__ DeviceReturnTy dhd() { return DeviceReturnTy(); }
+__host__ __device__ HostDeviceReturnTy dhd() { return HostDeviceReturnTy(); }
 
 // Same tests for extern "C" functions.
 extern "C" __host__ int chh() { return 0; } // expected-note {{previous definition is here}}
 extern "C" int chh() { return 0; }          // expected-error {{redefinition of 'chh'}}
 
 // H/D overloading is OK.
 extern "C" __device__ DeviceReturnTy cdh() { return DeviceReturnTy(); }
 extern "C" __host__ HostReturnTy cdh() { return HostReturnTy(); }
 
-// H/HD and D/HD overloading is not allowed.
-extern "C" __host__ __device__ int chhd1() { return 0; } // expected-note {{previous definition is here}}
-extern "C" __host__ int chhd1() { return 0; }            // expected-error {{redefinition of 'chhd1'}}
-
-extern "C" __host__ int chhd2() { return 0; }            // expected-note {{previous definition is here}}
-extern "C" __host__ __device__ int chhd2() { return 0; } // expected-error {{redefinition of 'chhd2'}}
-// expected-warning@-1 {{attribute declaration must precede definition}}
-// expected-note@-3 {{previous definition is here}}
+// H/HD and D/HD overloading is OK.
+extern "C" __host__ __device__ HostDeviceReturnTy chhd() { return HostDeviceReturnTy(); }
+extern "C" __host__ HostReturnTy chhd() { return HostReturnTy(); }
+
+extern "C" __host__ __device__ HostDeviceReturnTy chdd() { return HostDeviceReturnTy(); }
+extern "C" __device__ DeviceReturnTy chdd() { return DeviceReturnTy(); }
 
 // Helper functions to verify calling restrictions.
 __device__ DeviceReturnTy d() { return DeviceReturnTy(); }
 // expected-note@-1 0+ {{'d' declared here}}
 // expected-note@-2 0+ {{candidate function not viable: call to __device__ function from __host__ function}}
 // expected-note@-3 0+ {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
 
 __host__ HostReturnTy h() { return HostReturnTy(); }
__host__ void hostf() {
   HostFnPtr fp_ch = ch;
   HostReturnTy ret_ch = ch();
 
   HostFnPtr fp_dh = dh;
   HostReturnTy ret_dh = dh();
   HostFnPtr fp_cdh = cdh;
   HostReturnTy ret_cdh = cdh();
 
+  HostFnPtr fp_hdh = hdh;
+  HostReturnTy ret_hdh = hdh();
+  HostFnPtr fp_chhd = chhd;
+  HostReturnTy ret_chhd = chhd();
+
+  HostDeviceFnPtr fp_hdd = hdd;
+  HostDeviceReturnTy ret_hdd = hdd();
+  HostDeviceFnPtr fp_chdd = chdd;
+  HostDeviceReturnTy ret_chdd = chdd();
+
   GlobalFnPtr fp_g = g;
   g(); // expected-error {{call to global function g not configured}}
   g<<<0, 0>>>();
 }
 
 __device__ void devicef() {
   DeviceFnPtr fp_d = d;
   DeviceReturnTy ret_d = d();
   DeviceFnPtr fp_cd = cd;
   DeviceReturnTy ret_cd = cd();
 
   HostFnPtr fp_h = h;         // expected-error {{reference to __host__ function 'h' in __device__ function}}
   HostReturnTy ret_h = h();   // expected-error {{no matching function for call to 'h'}}
   HostFnPtr fp_ch = ch;       // expected-error {{reference to __host__ function 'ch' in __device__ function}}
   HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
 
   DeviceFnPtr fp_dh = dh;
   DeviceReturnTy ret_dh = dh();
   DeviceFnPtr fp_cdh = cdh;
   DeviceReturnTy ret_cdh = cdh();
 
+  HostDeviceFnPtr fp_hdh = hdh;
+  HostDeviceReturnTy ret_hdh = hdh();
+  HostDeviceFnPtr fp_chhd = chhd;
+  HostDeviceReturnTy ret_chhd = chhd();
+
+  DeviceFnPtr fp_hdd = hdd;
+  DeviceReturnTy ret_hdd = hdd();
+  DeviceFnPtr fp_chdd = chdd;
+  DeviceReturnTy ret_chdd = chdd();
+
   GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
   g(); // expected-error {{no matching function for call to 'g'}}
   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
 }
 
 __global__ void globalf() {
   DeviceFnPtr fp_d = d;
   DeviceReturnTy ret_d = d();
   DeviceFnPtr fp_cd = cd;
   DeviceReturnTy ret_cd = cd();
 
   HostFnPtr fp_h = h;         // expected-error {{reference to __host__ function 'h' in __global__ function}}
   HostReturnTy ret_h = h();   // expected-error {{no matching function for call to 'h'}}
   HostFnPtr fp_ch = ch;       // expected-error {{reference to __host__ function 'ch' in __global__ function}}
   HostReturnTy ret_ch = ch(); // expected-error {{no matching function for call to 'ch'}}
 
   DeviceFnPtr fp_dh = dh;
   DeviceReturnTy ret_dh = dh();
   DeviceFnPtr fp_cdh = cdh;
   DeviceReturnTy ret_cdh = cdh();
 
+  HostDeviceFnPtr fp_hdh = hdh;
+  HostDeviceReturnTy ret_hdh = hdh();
+  HostDeviceFnPtr fp_chhd = chhd;
+  HostDeviceReturnTy ret_chhd = chhd();
+
+  DeviceFnPtr fp_hdd = hdd;
+  DeviceReturnTy ret_hdd = hdd();
+  DeviceFnPtr fp_chdd = chdd;
+  DeviceReturnTy ret_chdd = chdd();
+
   GlobalFnPtr fp_g = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
   g(); // expected-error {{no matching function for call to 'g'}}
   g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
 }
 
 __host__ __device__ void hostdevicef() {
   DeviceFnPtr fp_d = d;
   DeviceReturnTy ret_d = d();
#if !defined(NOCHECKS) && defined(__CUDA_ARCH__)
   // expected-error@-5 {{no matching function for call to 'ch'}}
 #endif
 
   CurrentFnPtr fp_dh = dh;
   CurrentReturnTy ret_dh = dh();
   CurrentFnPtr fp_cdh = cdh;
   CurrentReturnTy ret_cdh = cdh();
 
+  // HDOrHostFoo is HostFoo if we're doing host compilation, and HDFoo
+  // otherwise.
+#ifdef __CUDA_ARCH__
+  typedef HostDeviceReturnTy HDOrHostReturnTy;
+  typedef HostDeviceFnPtr HDOrHostFnPtr;
+  typedef DeviceReturnTy HDOrDeviceReturnTy;
+  typedef DeviceFnPtr HDOrDeviceFnPtr;
+#else
+  typedef HostReturnTy HDOrHostReturnTy;
+  typedef HostFnPtr HDOrHostFnPtr;
+  typedef HostDeviceReturnTy HDOrDeviceReturnTy;
+  typedef HostDeviceFnPtr HDOrDeviceFnPtr;
+#endif
+
+  HDOrHostFnPtr fp_hdh = hdh;
+  HDOrHostReturnTy ret_hdh = hdh();
+  HDOrHostFnPtr fp_chhd = chhd;
+  HDOrHostReturnTy ret_chhd = chhd();
+
+  HDOrDeviceFnPtr fp_hdd = hdd;
+  HDOrDeviceReturnTy ret_hdd = hdd();
+  HDOrDeviceFnPtr fp_chdd = chdd;
+  HDOrDeviceReturnTy ret_chdd = chdd();
+
   GlobalFnPtr fp_g = g;
 #if defined(__CUDA_ARCH__)
   // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
 #endif
   g();
   g<<<0,0>>>();
 #if !defined(__CUDA_ARCH__)
   // expected-error@-3 {{call to global function g not configured}}
struct d_dh {
   __host__ ~d_dh() {}
 };
 
 // HD is OK
 struct d_hd {
   __host__ __device__ ~d_hd() {}
 };
 
-// Mixing H/D and HD is not allowed.
+// Mixing H/D and HD is OK.
 struct d_dhhd {
   __device__ ~d_dhhd() {}
-  __host__ ~d_dhhd() {} // expected-note {{previous declaration is here}}
-  __host__ __device__ ~d_dhhd() {} // expected-error {{destructor cannot be redeclared}}
+  __host__ ~d_dhhd() {}
+  __host__ __device__ ~d_dhhd() {}
 };
 
 struct d_hhd {
-  __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
-  __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
+  __host__ ~d_hhd() {}
+  __host__ __device__ ~d_hhd() {}
 };
 
 struct d_hdh {
-  __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
-  __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
+  __host__ __device__ ~d_hdh() {}
+  __host__ ~d_hdh() {}
 };
 
 struct d_dhd {
-  __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
-  __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
+  __device__ ~d_dhd() {}
+  __host__ __device__ ~d_dhd() {}
 };
 
 struct d_hdd {
-  __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
-  __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
+  __host__ __device__ ~d_hdd() {}
+  __device__ ~d_hdd() {}
 };
 
 // Test overloading of member functions
 struct m_h {
   void operator delete(void *ptr); // expected-note {{previous declaration is here}}
   __host__ void operator delete(void *ptr); // expected-error {{class member cannot be redeclared}}
 };
 
 // D/H overloading is OK
 struct m_dh {
   __device__ void operator delete(void *ptr);
   __host__ void operator delete(void *ptr);
 };
 
 // HD by itself is OK
 struct m_hd {
   __device__ __host__ void operator delete(void *ptr);
 };
 
 struct m_hhd {
-  __host__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
-  __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+  __host__ void operator delete(void *ptr) {}
+  __host__ __device__ void operator delete(void *ptr) {}
 };
 
 struct m_hdh {
-  __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
-  __host__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+  __host__ __device__ void operator delete(void *ptr) {}
+  __host__ void operator delete(void *ptr) {}
 };
 
 struct m_dhd {
-  __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
-  __host__ __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+  __device__ void operator delete(void *ptr) {}
+  __host__ __device__ void operator delete(void *ptr) {}
 };
 
 struct m_hdd {
-  __host__ __device__ void operator delete(void *ptr) {} // expected-note {{previous declaration is here}}
-  __device__ void operator delete(void *ptr) {} // expected-error {{class member cannot be redeclared}}
+  __host__ __device__ void operator delete(void *ptr) {}
+  __device__ void operator delete(void *ptr) {}
 };
 
 // __global__ functions can't be overloaded based on attribute
 // difference.
 struct G {
   friend void friend_of_g(G &arg);
 private:
   int x;

test/SemaCUDA/relaxed-constexpr.cu

+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s -fcuda-target-overloads -fcuda-relaxed-constexpr
+// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s -fcuda-target-overloads -fcuda-relaxed-constexpr -fcuda-is-device
+
+// expected-no-diagnostics
+
+#include "Inputs/cuda.h"
+
+static __device__ void f1();
+constexpr void f1();
+
+__device__ void f2();
+static constexpr void f2();
+
+// Different potential error depending on the order of declaration.
+constexpr void f3();
+static __device__ void f3();
+
+static constexpr void f4();
+__device__ void f4();
+
+// Variadic device functions are not allowed, so this is just treated as
+// host-only.
+constexpr void variadic(const char*, ...);