Add and use isDiscardableGVALinkage function.

[rnk]

lgtm

clang/include/clang/Sema/Sema.h

9244

I think you meant to make this change as part of the original patch.

[jlebar]

clang/include/clang/Sema/Sema.h

9244

Oops, thanks.

clang/include/clang/Basic/Linkage.h

 enum GVALinkage {
   GVA_Internal,
   GVA_AvailableExternally,
   GVA_DiscardableODR,
   GVA_StrongExternal,
   GVA_StrongODR
 };
 
+inline bool isDiscardableGVALinkage(GVALinkage L) {
+  return L <= GVA_DiscardableODR;
+}
+
 inline bool isExternallyVisible(Linkage L) {
   return L == ExternalLinkage || L == VisibleNoLinkage;
 }
 
 inline Linkage getFormalLinkage(Linkage L) {
   if (L == UniqueExternalLinkage)
     return ExternalLinkage;
   if (L == VisibleNoLinkage)

clang/include/clang/Sema/Sema.h

   llvm::DenseSet<FunctionDecl *> CUDAKnownEmittedFns;
 
   /// A partial call graph maintained during CUDA compilation to support
   /// deferred diagnostics.  Specifically, functions are only added here if, at
   /// the time they're added, they are not known-emitted.  As soon as we
   /// discover that a function is known-emitted, we remove it and everything it
   /// transitively calls from this set and add those functions to
   /// CUDAKnownEmittedFns.
-  llvm::DenseMap<FunctionDecl *, llvm::DenseSet<FunctionDecl *>> CUDACallGraph;
+  llvm::DenseMap<FunctionDecl *, llvm::SetVector<FunctionDecl *>> CUDACallGraph;
 
   /// Diagnostic builder for CUDA errors which may or may not be deferred.
   ///
   /// In CUDA, there exist constructs (e.g. variable-length arrays, try/catch)
   /// which are not allowed to appear inside __device__ functions and are
   /// allowed to appear in __host__ __device__ functions only if the host+device
   /// function is never codegen'ed.
   ///

clang/lib/AST/ASTContext.cpp

if (getTargetInfo().getCXXABI().canKeyFunctionBeInline()) {
         if (MD->isOutOfLine() && RD->isDynamicClass()) {
           const CXXMethodDecl *KeyFunc = getCurrentKeyFunction(RD);
           if (KeyFunc && KeyFunc->getCanonicalDecl() == MD->getCanonicalDecl())
             return true;
         }
       }
     }
 
-    GVALinkage Linkage = GetGVALinkageForFunction(FD);
-
     // static, static inline, always_inline, and extern inline functions can
     // always be deferred.  Normal inline functions can be deferred in C99/C++.
     // Implicit template instantiations can also be deferred in C++.
-    if (Linkage == GVA_Internal || Linkage == GVA_AvailableExternally ||
-        Linkage == GVA_DiscardableODR)
-      return false;
-    return true;
+    return !isDiscardableGVALinkage(GetGVALinkageForFunction(FD));
   }
   
   const VarDecl *VD = cast<VarDecl>(D);
   assert(VD->isFileVarDecl() && "Expected file scoped var");
 
   if (VD->isThisDeclarationADefinition() == VarDecl::DeclarationOnly &&
       !isMSStaticDataMemberInlineDefinition(VD))
     return false;
 
   // Variables that can be needed in other TUs are required.
-  GVALinkage L = GetGVALinkageForVariable(VD);
-  if (L != GVA_Internal && L != GVA_AvailableExternally &&
-      L != GVA_DiscardableODR)
+  if (!isDiscardableGVALinkage(GetGVALinkageForVariable(VD)))
     return true;
 
   // Variables that have destruction with side-effects are required.
   if (VD->getType().isDestructedType())
     return true;
 
   // Variables that have initialization with side-effects are required.
   if (VD->getInit() && VD->getInit()->HasSideEffects(*this) &&

clang/lib/Sema/SemaCUDA.cpp

static bool IsKnownEmitted(Sema &S, FunctionDecl *FD) {
   Sema::CUDAFunctionTarget T = S.IdentifyCUDATarget(FD);
   if (S.getLangOpts().CUDAIsDevice && T == Sema::CFT_Host)
     return false;
   if (!S.getLangOpts().CUDAIsDevice &&
       (T == Sema::CFT_Device || T == Sema::CFT_Global))
     return false;
 
   // Externally-visible and similar functions are always emitted.
-  if (S.getASTContext().GetGVALinkageForFunction(FD) > GVA_DiscardableODR)
+  if (!isDiscardableGVALinkage(S.getASTContext().GetGVALinkageForFunction(FD)))
     return true;
 
   // Otherwise, the function is known-emitted if it's in our set of
   // known-emitted functions.
   return S.CUDAKnownEmittedFns.count(FD) > 0;
 }
 
 Sema::CUDADiagBuilder Sema::CUDADiagIfDeviceCode(SourceLocation Loc,