[CUDA][HIP] Externalize kernels with internal linkage

This patch is a continuation of https://reviews.llvm.org/D123353.

Not only kernels in anonymous namespace, but also template
kernels with template arguments in anonymous namespace
need to be externalized.

To be more generic, this patch checks the linkage of a kernel
assuming the kernel does not have __global__ attribute. If
the linkage is internal then clang will externalize it.

This patch also fixes the postfix for externalized symbol
since nvptx does not allow '.' in symbol name.

Reviewed by: Artem Belevich

Differential Revision: https://reviews.llvm.org/D124189

Fixes: llvm#54560
(cherry picked from commit 04fb816)

clang/lib/CodeGen/CodeGenModule.cpp

@@ -1580,11 +1580,7 @@ StringRef CodeGenModule::getMangledName(GlobalDecl GD) {
   // directly between host- and device-compilations, the host- and
   // device-mangling in host compilation could help catching certain ones.
   assert(!isa<FunctionDecl>(ND) || !ND->hasAttr<CUDAGlobalAttr>() ||
-<<<<<<< HEAD
          getContext().shouldExternalize(ND) || getLangOpts().CUDAIsDevice ||
-=======
-         getContext().shouldExternalizeStaticVar(ND) || getLangOpts().CUDAIsDevice ||
->>>>>>> e6de9ed37308 ([CUDA][HIP] Externalize kernels in anonymous name space)
          (getContext().getAuxTargetInfo() &&
           (getContext().getAuxTargetInfo()->getCXXABI() !=
            getContext().getTargetInfo().getCXXABI())) ||
@@ -7153,7 +7149,6 @@ bool CodeGenModule::stopAutoInit() {
 
 void CodeGenModule::printPostfixForExternalizedDecl(llvm::raw_ostream &OS,
                                                     const Decl *D) const {
-<<<<<<< HEAD
   // ptxas does not allow '.' in symbol names. On the other hand, HIP prefers
   // postfix beginning with '.' since the symbol name can be demangled.
   if (LangOpts.HIP)
@@ -7218,8 +7213,4 @@ void CodeGenModule::moveLazyEmissionStates(CodeGenModule *NewBuilder) {
   NewBuilder->EmittedDeferredDecls = std::move(EmittedDeferredDecls);
 
   NewBuilder->ABI->MangleCtx = std::move(ABI->MangleCtx);
-=======
-  OS << (isa<VarDecl>(D) ? "__static__" : ".anon.")
-     << getContext().getCUIDHash();
->>>>>>> e6de9ed37308 ([CUDA][HIP] Externalize kernels in anonymous name space)
 }