[CUDA] Allow function overloads based on host/device attributes.

[tra]

Summary

The patch makes it possible to parse CUDA files that contain host/device functions with identical signatures, but different attributes without having to physically split source into host-only and device-only parts.

This change is needed in order to parse CUDA header files that have a lot of name clashes with standard include files.

Gory details are in design doc here: https://goo.gl/EXnymm
Feel free to leave comments there or in this review thread.

This feature is controlled with CC1 option -fcuda-target-overloads and is disabled by default.

[eliben]

include/clang/Sema/Sema.h

8599 ↗

(On Diff #33483)

The comment should explain exactly how it depends on the option

lib/AST/ItaniumMangle.cpp

504 ↗

(On Diff #33483)

Why device before host?

lib/Sema/SemaCUDA.cpp

98 ↗

(On Diff #33483)

Why allow Caller to be nullptr but not Callee? Also, != nullptr?

lib/Sema/SemaExprCXX.cpp

2277 ↗

(On Diff #33483)

Nit: wrap the ++I in braces too, otherwise this looks really weird

[tra]

addressed eliben@'s comments.

include/clang/Sema/Sema.h

8599 ↗

(On Diff #33483)

Rephrased.

lib/AST/ItaniumMangle.cpp

504 ↗

(On Diff #33483)

Name mangling chapter in Itanium CXX ABI says:

In cases where multiple order-insensitive qualifiers are present, they should be ordered [...] with the 'U' qualifiers in alphabetical order by the vendor name

So I've ordered attributes alphabetically.

lib/Sema/SemaCUDA.cpp

98 ↗

(On Diff #33483)

Caller may be null in global context (e.g. we're assigning function address to a global variable). Callee can't be. I'll add an assert().

Also, != nullptr?

I'm not quite sure what you mean here. Are you suggesting replacing "Caller ?" with "Caller != nullptr ?"

[eliben]

The CUDA parts look very good. Someone else should approve the overloading-related logic

include/clang/Sema/Sema.h

8605 ↗

(On Diff #33654)

Document this function

[I realize the other CUDA functions are not documented here :-(, but hey this is new code so it should follow the rules]

lib/Sema/SemaCUDA.cpp

98 ↗

(On Diff #33654)

Yes, I believe this is the accepted style when comparing pointers for null-ness.

106 ↗

(On Diff #33654)

Not just global from global. global from device too, right? As for global from HD, the CUDA guide forbids it

lib/Sema/SemaChecking.cpp

529 ↗

(On Diff #33654)

Is this part related to this patch?

[tra]

Removed builtin-related changes(D12122). Will commit them separately.
Added more test cases.
Addressed eliben@'s comments.

[tra]

lib/Sema/SemaCUDA.cpp

106 ↗

(On Diff #33654)

I'll update the comment.

As for HD->G, nvcc happily compiles following code:

__global__ void kernel() {}
__host__ __device__ void foo() {
#if !defined(__CUDA_ARCH__)
  kernel<<<0,0>>>();
#endif
}

Nvcc does produce an error for HD->G call during device compilation (the error actually complains about D->G or G->G calling). This patch matches nvcc behavior.

lib/Sema/SemaChecking.cpp

529 ↗

(On Diff #33654)

It's part of D12122 which broke some of your team's tests and got rolled back.
It's a prerequisite for overloads to work (otherwise anything that uses a builtin would violate calling convention either during device or during host compilation) and it also needs to be hidden behind some option so it does not break your tests again.

I think I can commit it separately after the overload patch. Overloading will not work with builtins until then, but I don't think it's a big deal as there are no users yet.

[rsmith]

include/clang/Basic/Builtins.h

85 ↗

(On Diff #33654)

I would prefer the more verbose isTargetBuiltin or isTargetSpecificBuiltin -- I don't think it will be obvious at call sites what this does if we use this abbreviation in the public interface.

lib/Sema/SemaExprCXX.cpp

2259 ↗

(On Diff #33654)

I don't see any test coverage for this; please add some tests that declare usual deallocation functions with CUDA host/device attributes and check that they behave as expected.

2276 ↗

(On Diff #33654)

Use resize, not set_size.

lib/Sema/SemaOverload.cpp

10119–10143 ↗

(On Diff #33654)

Please factor out the common code shared by this and FindUsualDeallocationFunction.

[tra]

Added more test cases to cover constructor/destructor/new/delete.
Refactored code to remove unwanted CUDA functions from Lookup match results.
Disabled HD overloading of H or D destructors to guarantee single destructor for codegen. (I'm still investigating whether it makes sense to enforce the same restriction on all functions as pcc@ suggested in design doc).

[rsmith]

lib/AST/ItaniumMangle.cpp

504 ↗

(On Diff #34059)

If we have enable_if on a __host__ __device__ function, I think we should mangle as Ua6deviceUa9enable_ifI...EUa4host. (The relative order of enable_if attributes matters, but the order of them relative to host and device does not, so retaining alphabetical order seems best.)

lib/Sema/SemaCUDA.cpp

216 ↗

(On Diff #34059)

We don't need the overhead of std::function here. Use this instead:

template<typename T, typename FetchDeclFn>
static void Erase...(
  ...,
  FetchDeclFn FetchDecl) {

[tra]

Implemented Richard Smiths' suggestions:
Fixed attribute order in name mangling.
Replaced std::function with template argument.

[tra]

Implemented pcc@'s suggestion to disallow mixing HD and H/D functions.
Removed name mangling as it's no longer needed. Restriction on HD overloading guarantees that we'll emit only one viable function during particular side of compilation (H or HD on the host side, D or HD on device side).
Updated test cases to reflect reduced number of attribute combinations we need to deal with now.

cfe/trunk/include/clang/Basic/LangOptions.def

 LANGOPT(NativeHalfType    , 1, 0, "Native half type support")
 LANGOPT(HalfArgsAndReturns, 1, 0, "half args and returns")
 LANGOPT(CUDA              , 1, 0, "CUDA")
 LANGOPT(OpenMP            , 1, 0, "OpenMP support")
 LANGOPT(OpenMPUseTLS      , 1, 0, "Use TLS for threadprivates or runtime calls")
 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(CUDAUsesLibDevice , 1, 0, "Selectively link and internalize bitcode.")
 
 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")

cfe/trunk/include/clang/Driver/CC1Options.td

 def fcuda_allow_host_calls_from_host_device : Flag<["-"],
     "fcuda-allow-host-calls-from-host-device">,
   HelpText<"Allow host device functions to call host functions">;
 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_uses_libdevice : Flag<["-"], "fcuda-uses-libdevice">,
   HelpText<"Selectively link and internalize bitcode.">;
 
 } // let Flags = [CC1Option]
 
 
 //===----------------------------------------------------------------------===//
 // cc1as-only Options

cfe/trunk/include/clang/Sema/Sema.h

enum CUDAFunctionTarget {
     CFT_Global,
     CFT_Host,
     CFT_HostDevice,
     CFT_InvalidTarget
   };
 
   CUDAFunctionTarget IdentifyCUDATarget(const FunctionDecl *D);
 
+  enum CUDAFunctionPreference {
+    CFP_Never,      // Invalid caller/callee combination.
+    CFP_LastResort, // Lowest priority. Only in effect if
+                    // LangOpts.CUDADisableTargetCallChecks is true.
+    CFP_Fallback,   // Low priority caller/callee combination
+    CFP_Best,       // Preferred caller/callee combination
+  };
+
+  /// Identifies relative preference of a given Caller/Callee
+  /// combination, based on their host/device attributes.
+  /// \param Caller function which needs address of \p Callee.
+  ///               nullptr in case of global context.
+  /// \param Callee target function
+  ///
+  /// \returns preference value for particular Caller/Callee combination.
+  CUDAFunctionPreference IdentifyCUDAPreference(const FunctionDecl *Caller,
+                                                const FunctionDecl *Callee);
+
   bool CheckCUDATarget(const FunctionDecl *Caller, const FunctionDecl *Callee);
 
+  /// Finds a function in \p Matches with highest calling priority
+  /// from \p Caller context and erases all functions with lower
+  /// calling priority.
+  void EraseUnwantedCUDAMatches(const FunctionDecl *Caller,
+                                SmallVectorImpl<FunctionDecl *> &Matches);
+  void EraseUnwantedCUDAMatches(const FunctionDecl *Caller,
+                                SmallVectorImpl<DeclAccessPair> &Matches);
+  void EraseUnwantedCUDAMatches(
+      const FunctionDecl *Caller,
+      SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches);
+
   /// Given a implicit special member, infer its CUDA target from the
   /// calls it needs to make to underlying base/field special members.
   /// \param ClassDecl the class for which the member is being created.
   /// \param CSM the kind of special member.
   /// \param MemberDecl the special member itself.
   /// \param ConstRHS true if this is a copy operation with a const object on
   ///        its RHS.
   /// \param Diagnose true if this call should emit diagnostics.

cfe/trunk/lib/Frontend/CompilerInvocation.cpp

if (Args.hasArg(OPT_fcuda_uses_libdevice))
     Opts.CUDAUsesLibDevice = 1;
 
   if (Args.hasArg(OPT_fcuda_allow_host_calls_from_host_device))
     Opts.CUDAAllowHostCallsFromHostDevice = 1;
 
   if (Args.hasArg(OPT_fcuda_disable_target_call_checks))
     Opts.CUDADisableTargetCallChecks = 1;
 
+  if (Args.hasArg(OPT_fcuda_target_overloads))
+    Opts.CUDATargetOverloads = 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))

cfe/trunk/lib/Sema/SemaCUDA.cpp

if (D->hasAttr<CUDADeviceAttr>()) {
     // Some implicit declarations (like intrinsic functions) are not marked.
     // Set the most lenient target on them for maximal flexibility.
     return CFT_HostDevice;
   }
 
   return CFT_Host;
 }
 
+// * CUDA Call preference table
+//
+// F - from,
+// T - to
+// Ph - preference in host mode
+// Pd - preference in device mode
+// H  - handled in (x)
+// Preferences: b-best, f-fallback, l-last resort, n-never.
+//
+// | F  | T  | Ph | Pd |  H  |
+// |----+----+----+----+-----+
+// | d  | d  | b  | b  | (b) |
+// | d  | g  | n  | n  | (a) |
+// | d  | h  | l  | l  | (e) |
+// | d  | hd | f  | f  | (c) |
+// | g  | d  | b  | b  | (b) |
+// | g  | g  | n  | n  | (a) |
+// | g  | h  | l  | l  | (e) |
+// | g  | hd | f  | f  | (c) |
+// | h  | d  | l  | l  | (e) |
+// | h  | g  | b  | b  | (b) |
+// | h  | h  | b  | b  | (b) |
+// | h  | hd | f  | f  | (c) |
+// | hd | d  | l  | f  | (d) |
+// | hd | g  | f  | n  |(d/a)|
+// | hd | h  | f  | l  | (d) |
+// | hd | hd | b  | b  | (b) |
+
+Sema::CUDAFunctionPreference
+Sema::IdentifyCUDAPreference(const FunctionDecl *Caller,
+                             const FunctionDecl *Callee) {
+  assert(getLangOpts().CUDATargetOverloads &&
+         "Should not be called w/o enabled target overloads.");
+
+  assert(Callee && "Callee must be valid.");
+  CUDAFunctionTarget CalleeTarget = IdentifyCUDATarget(Callee);
+  CUDAFunctionTarget CallerTarget =
+      (Caller != nullptr) ? IdentifyCUDATarget(Caller) : Sema::CFT_Host;
+
+  // If one of the targets is invalid, the check always fails, no matter what
+  // the other target is.
+  if (CallerTarget == CFT_InvalidTarget || CalleeTarget == CFT_InvalidTarget)
+    return CFP_Never;
+
+  // (a) Can't call global from some contexts until we support CUDA's
+  // dynamic parallelism.
+  if (CalleeTarget == CFT_Global &&
+      (CallerTarget == CFT_Global || CallerTarget == CFT_Device ||
+       (CallerTarget == CFT_HostDevice && getLangOpts().CUDAIsDevice)))
+    return CFP_Never;
+
+  // (b) Best case scenarios
+  if (CalleeTarget == CallerTarget ||
+      (CallerTarget == CFT_Host && CalleeTarget == CFT_Global) ||
+      (CallerTarget == CFT_Global && CalleeTarget == CFT_Device))
+    return CFP_Best;
+
+  // (c) Calling HostDevice is OK as a fallback that works for everyone.
+  if (CalleeTarget == CFT_HostDevice)
+    return CFP_Fallback;
+
+  // Figure out what should be returned 'last resort' cases. Normally
+  // those would not be allowed, but we'll consider them if
+  // CUDADisableTargetCallChecks is true.
+  CUDAFunctionPreference QuestionableResult =
+      getLangOpts().CUDADisableTargetCallChecks ? CFP_LastResort : CFP_Never;
+
+  // (d) HostDevice behavior depends on compilation mode.
+  if (CallerTarget == CFT_HostDevice) {
+    // Calling a function that matches compilation mode is OK.
+    // Calling a function from the other side is frowned upon.
+    if (getLangOpts().CUDAIsDevice)
+      return CalleeTarget == CFT_Device ? CFP_Fallback : QuestionableResult;
+    else
+      return (CalleeTarget == CFT_Host || CalleeTarget == CFT_Global)
+                 ? CFP_Fallback
+                 : QuestionableResult;
+  }
+
+  // (e) Calling across device/host boundary is not something you should do.
+  if ((CallerTarget == CFT_Host && CalleeTarget == CFT_Device) ||
+      (CallerTarget == CFT_Device && CalleeTarget == CFT_Host) ||
+      (CallerTarget == CFT_Global && CalleeTarget == CFT_Host))
+    return QuestionableResult;
+
+  llvm_unreachable("All cases should've been handled by now.");
+}
+
 bool Sema::CheckCUDATarget(const FunctionDecl *Caller,
                            const FunctionDecl *Callee) {
+  // With target overloads enabled, we only disallow calling
+  // combinations with CFP_Never.
+  if (getLangOpts().CUDATargetOverloads)
+    return IdentifyCUDAPreference(Caller,Callee) == CFP_Never;
+
   // The CUDADisableTargetCallChecks short-circuits this check: we assume all
   // cross-target calls are valid.
   if (getLangOpts().CUDADisableTargetCallChecks)
     return false;
 
   CUDAFunctionTarget CallerTarget = IdentifyCUDATarget(Caller),
                      CalleeTarget = IdentifyCUDATarget(Callee);
 
if (InDeviceMode && CalleeTarget != CFT_Device) {
 
       return true;
     }
   }
 
   return false;
 }
 
+template <typename T, typename FetchDeclFn>
+static void EraseUnwantedCUDAMatchesImpl(Sema &S, const FunctionDecl *Caller,
+                                         llvm::SmallVectorImpl<T> &Matches,
+                                         FetchDeclFn FetchDecl) {
+  assert(S.getLangOpts().CUDATargetOverloads &&
+         "Should not be called w/o enabled target overloads.");
+  if (Matches.size() <= 1)
+    return;
+
+  // Find the best call preference among the functions in Matches.
+  Sema::CUDAFunctionPreference P, BestCFP = Sema::CFP_Never;
+  for (auto const &Match : Matches) {
+    P = S.IdentifyCUDAPreference(Caller, FetchDecl(Match));
+    if (P > BestCFP)
+      BestCFP = P;
+  }
+
+  // Erase all functions with lower priority.
+  for (unsigned I = 0, N = Matches.size(); I != N;)
+    if (S.IdentifyCUDAPreference(Caller, FetchDecl(Matches[I])) < BestCFP) {
+      Matches[I] = Matches[--N];
+      Matches.resize(N);
+    } else {
+      ++I;
+    }
+}
+
+void Sema::EraseUnwantedCUDAMatches(const FunctionDecl *Caller,
+                                    SmallVectorImpl<FunctionDecl *> &Matches){
+  EraseUnwantedCUDAMatchesImpl<FunctionDecl *>(
+      *this, Caller, Matches, [](const FunctionDecl *item) { return item; });
+}
+
+void Sema::EraseUnwantedCUDAMatches(const FunctionDecl *Caller,
+                                    SmallVectorImpl<DeclAccessPair> &Matches) {
+  EraseUnwantedCUDAMatchesImpl<DeclAccessPair>(
+      *this, Caller, Matches, [](const DeclAccessPair &item) {
+        return dyn_cast<FunctionDecl>(item.getDecl());
+      });
+}
+
+void Sema::EraseUnwantedCUDAMatches(
+    const FunctionDecl *Caller,
+    SmallVectorImpl<std::pair<DeclAccessPair, FunctionDecl *>> &Matches){
+  EraseUnwantedCUDAMatchesImpl<std::pair<DeclAccessPair, FunctionDecl *>>(
+      *this, Caller, Matches,
+      [](const std::pair<DeclAccessPair, FunctionDecl *> &item) {
+        return dyn_cast<FunctionDecl>(item.second);
+      });
+}
+
 /// When an implicitly-declared special member has to invoke more than one
 /// base/field special member, conflicts may occur in the targets of these
 /// members. For example, if one base's member __host__ and another's is
 /// __device__, it's a conflict.
 /// This function figures out if the given targets \param Target1 and
 /// \param Target2 conflict, and if they do not it fills in
 /// \param ResolvedTarget with a target that resolves for both calls.
 /// \return true if there's a conflict, false otherwise.

cfe/trunk/lib/Sema/SemaDecl.cpp

 ///
 /// FIXME: This is a hack.
 template<typename T>
 static bool isIncompleteDeclExternC(Sema &S, const T *D) {
   if (S.getLangOpts().CPlusPlus) {
     // In C++, the overloadable attribute negates the effects of extern "C".
     if (!D->isInExternCContext() || D->template hasAttr<OverloadableAttr>())
       return false;
+
+    // So do CUDA's host/device attributes if overloading is enabled.
+    if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads &&
+        (D->template hasAttr<CUDADeviceAttr>() ||
+         D->template hasAttr<CUDAHostAttr>()))
+      return false;
   }
   return D->isExternC();
 }
 
 static bool shouldConsiderLinkage(const VarDecl *VD) {
   const DeclContext *DC = VD->getDeclContext()->getRedeclContext();
   if (DC->isFunctionOrMethod())
     return VD->hasExternalStorage();

cfe/trunk/lib/Sema/SemaExprCXX.cpp

if (getLangOpts().SizedDeallocation && Matches.size() == 2) {
     if (Matches[0]->getNumParams() != NumArgs)
       Matches.erase(Matches.begin());
     else
       Matches.erase(Matches.begin() + 1);
     assert(Matches[0]->getNumParams() == NumArgs &&
            "found an unexpected usual deallocation function");
   }
 
+  if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads)
+    EraseUnwantedCUDAMatches(dyn_cast<FunctionDecl>(CurContext), Matches);
+
   assert(Matches.size() == 1 &&
          "unexpectedly have multiple usual deallocation functions");
   return Matches.front();
 }
 
 bool Sema::FindDeallocationFunction(SourceLocation StartLoc, CXXRecordDecl *RD,
                                     DeclarationName Name,
                                     FunctionDecl* &Operator, bool Diagnose) {
for (LookupResult::iterator F = Found.begin(), FEnd = Found.end();
     // deallocation function.
     if (isa<FunctionTemplateDecl>(ND))
       continue;
 
     if (cast<CXXMethodDecl>(ND)->isUsualDeallocationFunction())
       Matches.push_back(F.getPair());
   }
 
+  if (getLangOpts().CUDA && getLangOpts().CUDATargetOverloads)
+    EraseUnwantedCUDAMatches(dyn_cast<FunctionDecl>(CurContext), Matches);
+
   // There's exactly one suitable operator;  pick it.
   if (Matches.size() == 1) {
     Operator = cast<CXXMethodDecl>(Matches[0]->getUnderlyingDecl());
 
     if (Operator->isDeleted()) {
       if (Diagnose) {
         Diag(StartLoc, diag::err_deleted_function_use);
         NoteDeletedFunction(Operator);

cfe/trunk/lib/Sema/SemaOverload.cpp

if (NewI == NewE || OldI == OldE)
       return true;
     llvm::FoldingSetNodeID NewID, OldID;
     NewI->getCond()->Profile(NewID, Context, true);
     OldI->getCond()->Profile(OldID, Context, true);
     if (NewID != OldID)
       return true;
   }
 
+  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 .
+    if ((NewTarget == CFT_HostDevice) || (OldTarget == CFT_HostDevice))
+      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.
   return false;
 }
 
 /// \brief Checks availability of the function depending on the current
 /// function context. Inside an unavailable function, unavailability is ignored.
 ///
 /// \returns true if \arg FD is unavailable and current context is inside
for (auto &Cand2A : Cand2Attrs) {
                                                      S.getASTContext(), true);
       if (Cand1ID != Cand2ID)
         return false;
     }
 
     return true;
   }
 
+  if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads &&
+      Cand1.Function && Cand2.Function) {
+    FunctionDecl *Caller = dyn_cast<FunctionDecl>(S.CurContext);
+    return S.IdentifyCUDAPreference(Caller, Cand1.Function) >
+           S.IdentifyCUDAPreference(Caller, Cand2.Function);
+  }
+
   return false;
 }
 
 /// \brief Computes the best viable function (C++ 13.3.3)
 /// within an overload candidate set.
 ///
 /// \param Loc The location of the function name (or operator symbol) for
 /// which overload resolution occurs.
if (FindAllFunctionsThatMatchTargetTypeExactly()) {
       //   If more than one function is selected, [...]
       if (Matches.size() > 1) {
         if (FoundNonTemplateFunction)
           EliminateAllTemplateMatches();
         else
           EliminateAllExceptMostSpecializedTemplate();
       }
     }
+
+    if (S.getLangOpts().CUDA && S.getLangOpts().CUDATargetOverloads &&
+        Matches.size() > 1)
+      EliminateSuboptimalCudaMatches();
   }
   
 private:
   bool isTargetTypeAFunction() const {
     return TargetFunctionType->isFunctionType();
   }
 
   // [ToType]     [Return]
private:
   void EliminateAllTemplateMatches() {
     //   [...] any function template specializations in the set are
     //   eliminated if the set also contains a non-template function, [...]
     for (unsigned I = 0, N = Matches.size(); I != N; ) {
       if (Matches[I].second->getPrimaryTemplate() == nullptr)
         ++I;
       else {
         Matches[I] = Matches[--N];
-        Matches.set_size(N);
-      }
-    }
+        Matches.resize(N);
+      }
+
+  void EliminateSuboptimalCudaMatches() {
+    S.EraseUnwantedCUDAMatches(dyn_cast<FunctionDecl>(S.CurContext), Matches);
   }
 
 public:
   void ComplainNoMatchesFound() const {
     assert(Matches.empty());
     S.Diag(OvlExpr->getLocStart(), diag::err_addr_ovl_no_viable)
         << OvlExpr->getName() << TargetFunctionType
         << OvlExpr->getSourceRange();

cfe/trunk/test/CodeGenCUDA/function-overload.cu

+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// Make sure we handle target overloads correctly.
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN:     -fcuda-target-overloads -emit-llvm -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device \
+// RUN:     -fcuda-target-overloads -emit-llvm -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE %s
+
+// Check target overloads handling with disabled call target checks.
+// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -emit-llvm \
+// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-HOST \
+// RUN:    -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-HOST-NC %s
+// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -emit-llvm \
+// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \
+// RUN:    -fcuda-is-device -o - %s \
+// RUN:     | FileCheck -check-prefix=CHECK-BOTH -check-prefix=CHECK-DEVICE \
+// RUN:    -check-prefix=CHECK-BOTH-NC -check-prefix=CHECK-DEVICE-NC %s
+
+#include "Inputs/cuda.h"
+
+typedef int (*fp_t)(void);
+typedef void (*gp_t)(void);
+
+// CHECK-HOST: @hp = global i32 ()* @_Z1hv
+// CHECK-HOST: @chp = global i32 ()* @ch
+// CHECK-HOST: @dhp = global i32 ()* @_Z2dhv
+// CHECK-HOST: @cdhp = global i32 ()* @cdh
+// CHECK-HOST: @gp = global void ()* @_Z1gv
+
+// CHECK-BOTH-LABEL: define i32 @_Z2dhv()
+__device__ int dh(void) { return 1; }
+// CHECK-DEVICE: ret i32 1
+__host__ int dh(void) { return 2; }
+// CHECK-HOST:   ret i32 2
+
+// CHECK-BOTH-LABEL: define i32 @_Z2hdv()
+__host__ __device__ int hd(void) { return 3; }
+// CHECK-BOTH:   ret i32 3
+
+// CHECK-DEVICE-LABEL: define i32 @_Z1dv()
+__device__ int d(void) { return 8; }
+// CHECK-DEVICE:   ret i32 8
+
+// CHECK-HOST-LABEL: define i32 @_Z1hv()
+__host__ int h(void) { return 9; }
+// CHECK-HOST:   ret i32 9
+
+// CHECK-BOTH-LABEL: define void @_Z1gv()
+__global__ void g(void) {}
+// CHECK-BOTH:   ret void
+
+// mangled names of extern "C" __host__ __device__ functions clash
+// with those of their __host__/__device__ counterparts, so
+// overloading of extern "C" functions can only happen for __host__
+// and __device__ functions -- we never codegen them in the same
+// compilation and therefore mangled name conflict is not a problem.
+
+// CHECK-BOTH-LABEL: define i32 @cdh()
+extern "C" __device__ int cdh(void) {return 10;}
+// CHECK-DEVICE:   ret i32 10
+extern "C" __host__ int cdh(void) {return 11;}
+// CHECK-HOST:     ret i32 11
+
+// CHECK-DEVICE-LABEL: define i32 @cd()
+extern "C" __device__ int cd(void) {return 12;}
+// CHECK-DEVICE:   ret i32 12
+
+// CHECK-HOST-LABEL: define i32 @ch()
+extern "C" __host__ int ch(void) {return 13;}
+// CHECK-HOST:     ret i32 13
+
+// CHECK-BOTH-LABEL: define i32 @chd()
+extern "C" __host__ __device__ int chd(void) {return 14;}
+// CHECK-BOTH:     ret i32 14
+
+// CHECK-HOST-LABEL: define void @_Z5hostfv()
+__host__ void hostf(void) {
+#if defined (NOCHECKS)
+  fp_t dp = d;   // CHECK-HOST-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
+  fp_t cdp = cd; // CHECK-HOST-NC: store {{.*}} @cd, {{.*}} %cdp,
+#endif
+  fp_t hp = h; // CHECK-HOST: store {{.*}} @_Z1hv, {{.*}} %hp,
+  fp_t chp = ch; // CHECK-HOST: store {{.*}} @ch, {{.*}} %chp,
+  fp_t dhp = dh; // CHECK-HOST: store {{.*}} @_Z2dhv, {{.*}} %dhp,
+  fp_t cdhp = cdh; // CHECK-HOST: store {{.*}} @cdh, {{.*}} %cdhp,
+  fp_t hdp = hd; // CHECK-HOST: store {{.*}} @_Z2hdv, {{.*}} %hdp,
+  fp_t chdp = chd; // CHECK-HOST: store {{.*}} @chd, {{.*}} %chdp,
+  gp_t gp = g; // CHECK-HOST: store {{.*}} @_Z1gv, {{.*}} %gp,
+
+#if defined (NOCHECKS)
+  d();     // CHECK-HOST-NC: call i32 @_Z1dv()
+  cd();    // CHECK-HOST-NC: call i32 @cd()
+#endif
+  h();     // CHECK-HOST: call i32 @_Z1hv()
+  ch();    // CHECK-HOST: call i32 @ch()
+  dh();    // CHECK-HOST: call i32 @_Z2dhv()
+  cdh();   // CHECK-HOST: call i32 @cdh()
+  g<<<0,0>>>();  // CHECK-HOST: call void @_Z1gv()
+}
+
+// CHECK-DEVICE-LABEL: define void @_Z7devicefv()
+__device__ void devicef(void) {
+  fp_t dp = d;   // CHECK-DEVICE: store {{.*}} @_Z1dv, {{.*}} %dp,
+  fp_t cdp = cd; // CHECK-DEVICE: store {{.*}} @cd, {{.*}} %cdp,
+#if defined (NOCHECKS)
+  fp_t hp = h; // CHECK-DEVICE-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
+  fp_t chp = ch; // CHECK-DEVICE-NC: store {{.*}} @ch, {{.*}} %chp,
+#endif
+  fp_t dhp = dh; // CHECK-DEVICE: store {{.*}} @_Z2dhv, {{.*}} %dhp,
+  fp_t cdhp = cdh; // CHECK-DEVICE: store {{.*}} @cdh, {{.*}} %cdhp,
+  fp_t hdp = hd; // CHECK-DEVICE: store {{.*}} @_Z2hdv, {{.*}} %hdp,
+  fp_t chdp = chd; // CHECK-DEVICE: store {{.*}} @chd, {{.*}} %chdp,
+
+  d();     // CHECK-DEVICE: call i32 @_Z1dv()
+  cd();    // CHECK-DEVICE: call i32 @cd()
+#if defined (NOCHECKS)
+  h();     // CHECK-DEVICE-NC: call i32 @_Z1hv()
+  ch();    // CHECK-DEVICE-NC: call i32 @ch()
+#endif
+  dh();    // CHECK-DEVICE: call i32 @_Z2dhv()
+  cdh();   // CHECK-DEVICE: call i32 @cdh()
+}
+
+// CHECK-BOTH-LABEL: define void @_Z11hostdevicefv()
+__host__ __device__ void hostdevicef(void) {
+#if defined (NOCHECKS)
+  fp_t dp = d;   // CHECK-BOTH-NC: store {{.*}} @_Z1dv, {{.*}} %dp,
+  fp_t cdp = cd; // CHECK-BOTH-NC: store {{.*}} @cd, {{.*}} %cdp,
+  fp_t hp = h; // CHECK-BOTH-NC: store {{.*}} @_Z1hv, {{.*}} %hp,
+  fp_t chp = ch; // CHECK-BOTH-NC: store {{.*}} @ch, {{.*}} %chp,
+#endif
+  fp_t dhp = dh; // CHECK-BOTH: store {{.*}} @_Z2dhv, {{.*}} %dhp,
+  fp_t cdhp = cdh; // CHECK-BOTH: store {{.*}} @cdh, {{.*}} %cdhp,
+  fp_t hdp = hd; // CHECK-BOTH: store {{.*}} @_Z2hdv, {{.*}} %hdp,
+  fp_t chdp = chd; // CHECK-BOTH: store {{.*}} @chd, {{.*}} %chdp,
+#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
+  gp_t gp = g; // CHECK-HOST-NC: store {{.*}} @_Z1gv, {{.*}} %gp,
+#endif
+
+#if defined (NOCHECKS)
+  d();     // CHECK-BOTH-NC: call i32 @_Z1dv()
+  cd();    // CHECK-BOTH-NC: call i32 @cd()
+  h();     // CHECK-BOTH-NC: call i32 @_Z1hv()
+  ch();    // CHECK-BOTH-NC: call i32 @ch()
+#endif
+  dh();    // CHECK-BOTH: call i32 @_Z2dhv()
+  cdh();   // CHECK-BOTH: call i32 @cdh()
+#if defined (NOCHECKS) && !defined(__CUDA_ARCH__)
+  g<<<0,0>>>();  // CHECK-HOST-NC: call void @_Z1gv()
+#endif
+}
+
+// Test for address of overloaded function resolution in the global context.
+fp_t hp = h;
+fp_t chp = ch;
+fp_t dhp = dh;
+fp_t cdhp = cdh;
+gp_t gp = g;
+
+int x;
+// Check constructors/destructors for D/H functions
+struct s_cd_dh {
+  __host__ s_cd_dh() { x = 11; }
+  __device__ s_cd_dh() { x = 12; }
+  __host__ ~s_cd_dh() { x = 21; }
+  __device__ ~s_cd_dh() { x = 22; }
+};
+
+struct s_cd_hd {
+  __host__ __device__ s_cd_hd() { x = 31; }
+  __host__ __device__ ~s_cd_hd() { x = 32; }
+};
+
+// CHECK-BOTH: define void @_Z7wrapperv
+#if defined(__CUDA_ARCH__)
+__device__
+#else
+__host__
+#endif
+void wrapper() {
+  s_cd_dh scddh;
+  // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(
+  s_cd_hd scdhd;
+  // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev
+
+  // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(
+  // CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev(
+}
+// CHECK-BOTH: ret void
+
+// Now it's time to check what's been generated for the methods we used.
+
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev(
+// CHECK-HOST:   store i32 11,
+// CHECK-DEVICE: store i32 12,
+// CHECK-BOTH: ret void
+
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
+// CHECK-BOTH:   store i32 31,
+// CHECK-BOTH: ret void
+
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
+// CHECK-BOTH: store i32 32,
+// CHECK-BOTH: ret void
+
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev(
+// CHECK-HOST:   store i32 21,
+// CHECK-DEVICE: store i32 22,
+// CHECK-BOTH: ret void
+

cfe/trunk/test/SemaCUDA/function-overload.cu

+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// Make sure we handle target overloads correctly.
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \
+// RUN:    -fsyntax-only -fcuda-target-overloads -verify %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda \
+// RUN:    -fsyntax-only -fcuda-target-overloads -fcuda-is-device -verify %s
+
+// Check target overloads handling with disabled call target checks.
+// RUN: %clang_cc1 -DNOCHECKS -triple x86_64-unknown-linux-gnu -fsyntax-only \
+// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads -verify %s
+// RUN: %clang_cc1 -DNOCHECKS -triple nvptx64-nvidia-cuda -fsyntax-only \
+// RUN:    -fcuda-disable-target-call-checks -fcuda-target-overloads \
+// RUN:    -fcuda-is-device -verify %s
+
+#include "Inputs/cuda.h"
+
+typedef int (*fp_t)(void);
+typedef void (*gp_t)(void);
+
+// Host and unattributed functions can't be overloaded
+__host__ int hh(void) { return 1; } // expected-note {{previous definition is here}}
+int hh(void) { return 1; } // expected-error {{redefinition of 'hh'}}
+
+// H/D overloading is OK
+__host__ int dh(void) { return 2; }
+__device__ int dh(void) { return 2; }
+
+// H/HD and D/HD are not allowed
+__host__ __device__ int hdh(void) { return 5; } // expected-note {{previous definition is here}}
+__host__ int hdh(void) { return 4; } // expected-error {{redefinition of 'hdh'}}
+
+__host__ int hhd(void) { return 4; } // expected-note {{previous definition is here}}
+__host__ __device__ int hhd(void) { return 5; } // expected-error {{redefinition of 'hhd'}}
+// expected-warning@-1 {{attribute declaration must precede definition}}
+// expected-note@-3 {{previous definition is here}}
+
+__host__ __device__ int hdd(void) { return 7; } // expected-note {{previous definition is here}}
+__device__ int hdd(void) { return 6; } // expected-error {{redefinition of 'hdd'}}
+
+__device__ int dhd(void) { return 6; } // expected-note {{previous definition is here}}
+__host__ __device__ int dhd(void) { return 7; } // expected-error {{redefinition of 'dhd'}}
+// expected-warning@-1 {{attribute declaration must precede definition}}
+// expected-note@-3 {{previous definition is here}}
+
+// Same tests for extern "C" functions
+extern "C" __host__ int chh(void) {return 11;} // expected-note {{previous definition is here}}
+extern "C" int chh(void) {return 11;} // expected-error {{redefinition of 'chh'}}
+
+// H/D overloading is OK
+extern "C" __device__ int cdh(void) {return 10;}
+extern "C" __host__ int cdh(void) {return 11;}
+
+// H/HD and D/HD overloading is not allowed.
+extern "C" __host__ __device__ int chhd1(void) {return 12;} // expected-note {{previous definition is here}}
+extern "C" __host__ int chhd1(void) {return 13;} // expected-error {{redefinition of 'chhd1'}}
+
+extern "C" __host__ int chhd2(void) {return 13;} // expected-note {{previous definition is here}}
+extern "C" __host__ __device__ int chhd2(void) {return 12;} // expected-error {{redefinition of 'chhd2'}}
+// expected-warning@-1 {{attribute declaration must precede definition}}
+// expected-note@-3 {{previous definition is here}}
+
+// Helper functions to verify calling restrictions.
+__device__ int d(void) { return 8; }
+__host__ int h(void) { return 9; }
+__global__ void g(void) {}
+extern "C" __device__ int cd(void) {return 10;}
+extern "C" __host__ int ch(void) {return 11;}
+
+__host__ void hostf(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{reference to __device__ function 'd' in __host__ function}}
+  // expected-note@65 {{'d' declared here}}
+  // expected-error@-4 {{reference to __device__ function 'cd' in __host__ function}}
+  // expected-note@68 {{'cd' declared here}}
+#endif
+  fp_t hp = h;
+  fp_t chp = ch;
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g;
+
+  d();
+  cd();
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{no matching function for call to 'd'}}
+  // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ function}}
+  // expected-error@-4 {{no matching function for call to 'cd'}}
+  // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ function}}
+#endif
+  h();
+  ch();
+  dh();
+  cdh();
+  g(); // expected-error {{call to global function g not configured}}
+  g<<<0,0>>>();
+}
+
+
+__device__ void devicef(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+  fp_t hp = h;
+  fp_t chp = ch;
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{reference to __host__ function 'h' in __device__ function}}
+  // expected-note@66 {{'h' declared here}}
+  // expected-error@-4 {{reference to __host__ function 'ch' in __device__ function}}
+  // expected-note@69 {{'ch' declared here}}
+#endif
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __device__ function}}
+               // expected-note@67 {{'g' declared here}}
+
+  d();
+  cd();
+  h();
+  ch();
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{no matching function for call to 'h'}}
+  // expected-note@66 {{candidate function not viable: call to __host__ function from __device__ function}}
+  // expected-error@-4 {{no matching function for call to 'ch'}}
+  // expected-note@69 {{candidate function not viable: call to __host__ function from __device__ function}}
+#endif
+  dh();
+  cdh();
+  g(); // expected-error {{no matching function for call to 'g'}}
+  // expected-note@67 {{candidate function not viable: call to __global__ function from __device__ function}}
+  g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __device__ function}}
+  // expected-note@67 {{'g' declared here}}
+}
+
+__global__ void globalf(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+  fp_t hp = h;
+  fp_t chp = ch;
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{reference to __host__ function 'h' in __global__ function}}
+  // expected-note@66 {{'h' declared here}}
+  // expected-error@-4 {{reference to __host__ function 'ch' in __global__ function}}
+  // expected-note@69 {{'ch' declared here}}
+#endif
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g; // expected-error {{reference to __global__ function 'g' in __global__ function}}
+               // expected-note@67 {{'g' declared here}}
+
+  d();
+  cd();
+  h();
+  ch();
+#if !defined(NOCHECKS)
+  // expected-error@-3 {{no matching function for call to 'h'}}
+  // expected-note@66 {{candidate function not viable: call to __host__ function from __global__ function}}
+  // expected-error@-4 {{no matching function for call to 'ch'}}
+  // expected-note@69 {{candidate function not viable: call to __host__ function from __global__ function}}
+#endif
+  dh();
+  cdh();
+  g(); // expected-error {{no matching function for call to 'g'}}
+  // expected-note@67 {{candidate function not viable: call to __global__ function from __global__ function}}
+  g<<<0,0>>>(); // expected-error {{reference to __global__ function 'g' in __global__ function}}
+  // expected-note@67 {{'g' declared here}}
+}
+
+__host__ __device__ void hostdevicef(void) {
+  fp_t dp = d;
+  fp_t cdp = cd;
+  fp_t hp = h;
+  fp_t chp = ch;
+#if !defined(NOCHECKS)
+#if !defined(__CUDA_ARCH__)
+  // expected-error@-6 {{reference to __device__ function 'd' in __host__ __device__ function}}
+  // expected-note@65 {{'d' declared here}}
+  // expected-error@-7 {{reference to __device__ function 'cd' in __host__ __device__ function}}
+  // expected-note@68 {{'cd' declared here}}
+#else
+  // expected-error@-9 {{reference to __host__ function 'h' in __host__ __device__ function}}
+  // expected-note@66 {{'h' declared here}}
+  // expected-error@-10 {{reference to __host__ function 'ch' in __host__ __device__ function}}
+  // expected-note@69 {{'ch' declared here}}
+#endif
+#endif
+  fp_t dhp = dh;
+  fp_t cdhp = cdh;
+  gp_t gp = g;
+#if defined(__CUDA_ARCH__)
+  // expected-error@-2 {{reference to __global__ function 'g' in __host__ __device__ function}}
+  // expected-note@67 {{'g' declared here}}
+#endif
+
+  d();
+  cd();
+  h();
+  ch();
+#if !defined(NOCHECKS)
+#if !defined(__CUDA_ARCH__)
+  // expected-error@-6 {{no matching function for call to 'd'}}
+  // expected-note@65 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+  // expected-error@-7 {{no matching function for call to 'cd'}}
+  // expected-note@68 {{candidate function not viable: call to __device__ function from __host__ __device__ function}}
+#else
+  // expected-error@-9 {{no matching function for call to 'h'}}
+  // expected-note@66 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+  // expected-error@-10 {{no matching function for call to 'ch'}}
+  // expected-note@69 {{candidate function not viable: call to __host__ function from __host__ __device__ function}}
+#endif
+#endif
+
+  dh();
+  cdh();
+  g();
+  g<<<0,0>>>();
+#if !defined(__CUDA_ARCH__)
+  // expected-error@-3 {{call to global function g not configured}}
+#else
+  // expected-error@-5 {{no matching function for call to 'g'}}
+  // expected-note@67 {{candidate function not viable: call to __global__ function from __host__ __device__ function}}
+  // expected-error@-6 {{reference to __global__ function 'g' in __host__ __device__ function}}
+  // expected-note@67 {{'g' declared here}}
+#endif  // __CUDA_ARCH__
+}
+
+// Test for address of overloaded function resolution in the global context.
+fp_t hp = h;
+fp_t chp = ch;
+fp_t dhp = dh;
+fp_t cdhp = cdh;
+gp_t gp = g;
+
+
+// Test overloading of destructors
+// Can't mix H and unattributed destructors
+struct d_h {
+  ~d_h() {} // expected-note {{previous declaration is here}}
+  __host__ ~d_h() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+// H/D overloading is OK
+struct d_dh {
+  __device__ ~d_dh() {}
+  __host__ ~d_dh() {}
+};
+
+// HD is OK
+struct d_hd {
+  __host__ __device__ ~d_hd() {}
+};
+
+// Mixing H/D and HD is not allowed.
+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}}
+};
+
+struct d_hhd {
+  __host__ ~d_hhd() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~d_hhd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_hdh {
+  __host__ __device__ ~d_hdh() {} // expected-note {{previous declaration is here}}
+  __host__ ~d_hdh() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_dhd {
+  __device__ ~d_dhd() {} // expected-note {{previous declaration is here}}
+  __host__ __device__ ~d_dhd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+struct d_hdd {
+  __host__ __device__ ~d_hdd() {} // expected-note {{previous declaration is here}}
+  __device__ ~d_hdd() {} // expected-error {{destructor cannot be redeclared}}
+};
+
+// 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}}
+};
+
+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}}
+};
+
+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}}
+};
+
+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}}
+};