[CUDA][HIP] Externalize kernels with internal linkage

[yaxunl]

Summary

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.

[tra]

clang/lib/AST/ASTContext.cpp

11323

Nit: I'd phrase it as a positive assertion ConsiderCudaGlobalAttr and default it to true.

DontDoX always trips me and gets me to question it -- "what *are* we doing then? what else is there besides X?".
With a DoX things are usually simpler and limited to X -- we're either doing X or not.

12307

Perhaps we don't need to change the public AST API and plumb IgnoreCUDAGlobalAttr through.
We cold create CUDA-aware static version of GetGVALinkageForCudaKernel instead, which would call adjustGVALinkageForExternalDefinitionKind(..., adjustGVALinkageForAttributes(IgnoreCUDAGlobalAttr=true)).

[yaxunl]

clang/lib/AST/ASTContext.cpp

11323

will do

12307

We could have a static function but it would be GetGVALinkageForCUDAKernelWithoutGlobalAttr since we need to know the linkage of the kernel assuming it has no __global__ attribute.

If you think it is OK I can make the change.

[yaxunl]

use static function

[tra]

clang/lib/AST/ASTContext.cpp

12307

No point making public what's of no use to anybody other than this particular instance.

To think of it, we don't even need a function and could just do

if (D->hasAttr<CUDAGlobalAttr>() ) {
  bool OriginalKernelLinkage = adjustGVALinkageForExternalDefinitionKind(..., adjustGVALinkageForAttributes(IgnoreCUDAGlobalAttr=true));
  return OriginalKernelLinkage == GVA_Internal;
}
return (IsStaticVar &&....)

[yaxunl]

clang/lib/AST/ASTContext.cpp

12307

One disadvantage of this approach is that it duplicates the code in GetGVALinkageForFunction. In the future, if GetGVALinkageForFunction changes, the same change needs to be applied to the duplicated code, which is error-prone.

[tra]

clang/lib/AST/ASTContext.cpp

12307

Good point. Looking at the code closer, it appears that what we're interested in is whether the kernel was internal and *became* externally visible due to it being a kernel.

Right now we're checking if the function would normally be GVA_Internal (shouldn't we have considered GVA_DiscardableODR, too, BTW?)
This is a somewhat indirect way of figuring out what we really need.

The code that determines what we want is essentially this code in adjustGVALinkageForAttributes that we're trying to enable/disable with ConsiderCudaGlobalAttr.

It can be easily extracted into a static function, which could then be used from both adjustGVALinkageForAttributes, (which would no longer need ConsiderCudaGlobalAttr) and from here.

bool isInternalKernel(ASTContext *Context, Decl *D) {
  L=basicGVALinkageForFunction(Context, D);
  return (D->hasAttr<CUDAGlobalAttr>() &&
          (L == GVA_DiscardableODR || L == GVA_Internal));
}

This would both avoid logic duplication and would better match our intent.

Does it make sense? Or did I miss something else?

[yaxunl]

clang/lib/AST/ASTContext.cpp

12307

GVA_DiscardableODR usually maps to linkonce_odr linkage in LLVM IR. It follows the ODR, therefore we should not make them unique.

If we use isInternalKernel in adjustGVALinkageForAttributes, there will be two calls of basicGVALinkageForFunction when GetGVALinkageForFunction is called, which seems inefficient. I think we can keep GetGVALinkageForFunction as it was, and use basicGVALinkageForFunction directly in mayExternalize.

[yaxunl]

use basicGVALinkageForFunction

[tra]

LGTM overal, with few test nits.

clang/lib/AST/ASTContext.cpp

12307

SGTM.

clang/test/CodeGenCUDA/device-var-linkage.cu

1–2

This is odd -- the tests use -x hip and -triple nvptx.

I think we need to change them into HIP+amdgpu and CUDA +nvptx variants ans we now have language-dependent behavior here and are interested in the language/triple combinations that we do use in practice.

clang/test/CodeGenCUDA/kernel-in-anon-ns.cu

3

We should have CUDA test variants here, too.

clang/test/CodeGenCUDA/managed-var.cu

1

Tests above do not have REQUIRED. Is it needed here?

[yaxunl]

clang/test/CodeGenCUDA/device-var-linkage.cu

1–2

will change them to amdgcn and add CUDA variant when committing.

clang/test/CodeGenCUDA/kernel-in-anon-ns.cu

3

will add CUDA test when committing.

clang/test/CodeGenCUDA/managed-var.cu

1

No. I will remove it when committing.

clang/lib/AST/ASTContext.cpp

   // forcibly get emitted.  While the body of the function cannot be later
   // replaced, the function definition cannot be discarded.
   if (FD->isMSExternInline())
     return GVA_StrongODR;
 
   return GVA_DiscardableODR;
 }
 
-static GVALinkage adjustGVALinkageForAttributes(const ASTContext &Context,
-                                                const Decl *D, GVALinkage L) {
+static GVALinkage
+adjustGVALinkageForAttributes(const ASTContext &Context, const Decl *D,
+                              GVALinkage L,
+                              bool ConsiderCudaGlobalAttr = true) {
   // See http://msdn.microsoft.com/en-us/library/xa0d9ste.aspx
   // dllexport/dllimport on inline functions.
   if (D->hasAttr<DLLImportAttr>()) {
     if (L == GVA_DiscardableODR || L == GVA_StrongODR)
       return GVA_AvailableExternally;
   } else if (D->hasAttr<DLLExportAttr>()) {
     if (L == GVA_DiscardableODR)
       return GVA_StrongODR;
-  } else if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice) {
+  } else if (Context.getLangOpts().CUDA && Context.getLangOpts().CUDAIsDevice &&
+             ConsiderCudaGlobalAttr) {
     // Device-side functions with __global__ attribute must always be
     // visible externally so they can be launched from host.
     if (D->hasAttr<CUDAGlobalAttr>() &&
         (L == GVA_DiscardableODR || L == GVA_Internal))
       return GVA_StrongODR;
     // Single source offloading languages like CUDA/HIP need to be able to
     // access static device variables from host code of the same compilation
     // unit. This is done by externalizing the static variable with a shared

 }
 
 GVALinkage ASTContext::GetGVALinkageForFunction(const FunctionDecl *FD) const {
   return adjustGVALinkageForExternalDefinitionKind(*this, FD,
            adjustGVALinkageForAttributes(*this, FD,
              basicGVALinkageForFunction(*this, FD)));
 }
 
+static GVALinkage
+GetGVALinkageForCUDAKernelWithoutGlobalAttr(const ASTContext &Context,
+                                            const FunctionDecl *FD) {
+  return adjustGVALinkageForExternalDefinitionKind(
+      Context, FD,
+      adjustGVALinkageForAttributes(Context, FD,
+                                    basicGVALinkageForFunction(Context, FD),
+                                    /*ConsiderCudaGlobalAttr=*/false));
+}
+
 static GVALinkage basicGVALinkageForVariable(const ASTContext &Context,
                                              const VarDecl *VD) {
   if (!VD->isExternallyVisible())
     return GVA_Internal;
 
   if (VD->isStaticLocal()) {
     const DeclContext *LexicalContext = VD->getParentFunctionOrMethod();
     while (LexicalContext && !isa<FunctionDecl>(LexicalContext))

                               !D->getAttr<CUDADeviceAttr>()->isImplicit()) ||
                              (D->hasAttr<CUDAConstantAttr>() &&
                               !D->getAttr<CUDAConstantAttr>()->isImplicit());
   // CUDA/HIP: static managed variables need to be externalized since it is
   // a declaration in IR, therefore cannot have internal linkage. Kernels in
   // anonymous name space needs to be externalized to avoid duplicate symbols.
   return (IsStaticVar &&
           (D->hasAttr<HIPManagedAttr>() || IsExplicitDeviceVar)) ||
-         (D->hasAttr<CUDAGlobalAttr>() && D->isInAnonymousNamespace());
+         (D->hasAttr<CUDAGlobalAttr>() &&
+          GetGVALinkageForCUDAKernelWithoutGlobalAttr(
+              *this, cast<FunctionDecl>(D)) == GVA_Internal);
 }
 
 bool ASTContext::shouldExternalize(const Decl *D) const {
   return mayExternalize(D) &&
          (D->hasAttr<HIPManagedAttr>() || D->hasAttr<CUDAGlobalAttr>() ||
           CUDADeviceVarODRUsedByHost.count(cast<VarDecl>(D)));
 }
 
 StringRef ASTContext::getCUIDHash() const {
   if (!CUIDHash.empty())
     return CUIDHash;
   if (LangOpts.CUID.empty())
     return StringRef();
   CUIDHash = llvm::utohexstr(llvm::MD5Hash(LangOpts.CUID), /*LowerCase=*/true);
   return CUIDHash;
 }

clang/lib/CodeGen/CodeGenModule.cpp

if (StopAfter) {
     }
     ++NumAutoVarInit;
   }
   return false;
 }
 
 void CodeGenModule::printPostfixForExternalizedDecl(llvm::raw_ostream &OS,
                                                     const Decl *D) const {
-  OS << (isa<VarDecl>(D) ? "__static__" : ".anon.")
-     << getContext().getCUIDHash();
+  StringRef Tag;
+  // 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)
+    Tag = (isa<VarDecl>(D) ? ".static." : ".intern.");
+  else
+    Tag = (isa<VarDecl>(D) ? "__static__" : "__intern__");
+  OS << Tag << getContext().getCUIDHash();
 }
 
 namespace {
 class NoLoopChecker final : public ConstStmtVisitor<NoLoopChecker> {
 public:
   NoLoopChecker() : RejectNoLoop{false} {}
   bool isRejectNoLoop() const { return RejectNoLoop; }
 

clang/test/CodeGenCUDA/device-var-linkage.cu

 // RUN: %clang_cc1 -no-opaque-pointers -triple nvptx -fcuda-is-device \
 // RUN:   -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=DEV,NORDC %s
 // RUN: %clang_cc1 -no-opaque-pointers -triple nvptx -fcuda-is-device \
 // RUN:   -fgpu-rdc -cuid=abc -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=DEV,RDC %s
 // RUN: %clang_cc1 -no-opaque-pointers -triple nvptx \
 // RUN:   -emit-llvm -o - -x hip %s \
 // RUN:   | FileCheck -check-prefixes=HOST,NORDC-H %s
 // RUN: %clang_cc1 -no-opaque-pointers -triple nvptx \

 // DEV-DAG: @ev2 = external addrspace(4) global i32
 // HOST-DAG: @ev2 = external global i32
 extern __constant__ int ev2;
 // DEV-DAG: @ev3 = external addrspace(1) externally_initialized global i32 addrspace(1)*
 // HOST-DAG: @ev3 = external externally_initialized global i32*
 extern __managed__ int ev3;
 
 // NORDC-DAG: @_ZL3sv1 = addrspace(1) externally_initialized global i32 0
-// RDC-DAG: @_ZL3sv1__static__[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
+// RDC-DAG: @_ZL3sv1.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
 // HOST-DAG: @_ZL3sv1 = internal global i32 undef
 static __device__ int sv1;
 // NORDC-DAG: @_ZL3sv2 = addrspace(4) externally_initialized global i32 0
-// RDC-DAG: @_ZL3sv2__static__[[HASH]] = addrspace(4) externally_initialized global i32 0
+// RDC-DAG: @_ZL3sv2.static.[[HASH]] = addrspace(4) externally_initialized global i32 0
 // HOST-DAG: @_ZL3sv2 = internal global i32 undef
 static __constant__ int sv2;
 // NORDC-DAG: @_ZL3sv3 = addrspace(1) externally_initialized global i32 addrspace(1)* null
-// RDC-DAG: @_ZL3sv3__static__[[HASH]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
+// RDC-DAG: @_ZL3sv3.static.[[HASH]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
 // HOST-DAG: @_ZL3sv3 = internal externally_initialized global i32* null
 static __managed__ int sv3;
 
 __device__ __host__ int work(int *x);
 
 __device__ __host__ int fun1() {
   return work(&ev1) + work(&ev2) + work(&ev3) + work(&sv1) + work(&sv2) + work(&sv3);
 }

clang/test/CodeGenCUDA/kernel-in-anon-ns.cu

 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
 // RUN:   -aux-triple x86_64-unknown-linux-gnu -std=c++11 -fgpu-rdc \
 // RUN:   -emit-llvm -o - -x hip %s > %t.dev
 
 // RUN: %clang_cc1 -triple x86_64-gnu-linux -cuid=abc \
 // RUN:   -aux-triple amdgcn-amd-amdhsa -std=c++11 -fgpu-rdc \
 // RUN:   -emit-llvm -o - -x hip %s > %t.host
 
 // RUN: cat %t.dev %t.host | FileCheck %s
 
+// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -cuid=abc \
+// RUN:   -aux-triple x86_64-unknown-linux-gnu -std=c++11 -fgpu-rdc \
+// RUN:   -emit-llvm -o - %s | FileCheck -check-prefix=CUDA %s
+
 #include "Inputs/cuda.h"
 
-// CHECK: define weak_odr {{.*}}void @[[KERN:_ZN12_GLOBAL__N_16kernelEv\.anon\.b04fd23c98500190]](
-// CHECK: @[[STR:.*]] = {{.*}} c"[[KERN]]\00"
-// CHECK: call i32 @__hipRegisterFunction({{.*}}@[[STR]]
+// CHECK-DAG: define weak_odr {{.*}}void @[[KERN1:_ZN12_GLOBAL__N_16kernelEv\.intern\.b04fd23c98500190]](
+// CHECK-DAG: define weak_odr {{.*}}void @[[KERN2:_Z8tempKernIN12_GLOBAL__N_11XEEvT_\.intern\.b04fd23c98500190]](
+// CHECK-DAG: define weak_odr {{.*}}void @[[KERN3:_Z8tempKernIN12_GLOBAL__N_1UlvE_EEvT_\.intern\.b04fd23c98500190]](
+// CHECK-DAG: @[[STR1:.*]] = {{.*}} c"[[KERN1]]\00"
+// CHECK-DAG: @[[STR2:.*]] = {{.*}} c"[[KERN2]]\00"
+// CHECK-DAG: @[[STR3:.*]] = {{.*}} c"[[KERN3]]\00"
+// CHECK-DAG: call i32 @__hipRegisterFunction({{.*}}@[[STR1]]
+// CHECK-DAG: call i32 @__hipRegisterFunction({{.*}}@[[STR2]]
+// CHECK-DAG: call i32 @__hipRegisterFunction({{.*}}@[[STR3]]
+
+// CUDA: define weak_odr {{.*}}void @[[KERN1:_ZN12_GLOBAL__N_16kernelEv__intern__b04fd23c98500190]](
+
+template <typename T>
+__global__ void tempKern(T x) {}
 
 namespace {
-__global__ void kernel() {
-}
+  __global__ void kernel() {}
+  struct X {};
+  X x;
+  auto lambda = [](){};
 }
 
 void test() {
   kernel<<<1, 1>>>();
+
+  tempKern<<<1, 1>>>(x);
+
+  tempKern<<<1, 1>>>(lambda);
 }

clang/test/CodeGenCUDA/managed-var.cu

 // REQUIRES: x86-registered-target, amdgpu-registered-target
 
 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
 // RUN:   -emit-llvm -o - -x hip %s | FileCheck \
 // RUN:   -check-prefixes=COMMON,DEV,NORDC-D %s
 
 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
 // RUN:   -emit-llvm -fgpu-rdc -cuid=abc -o - -x hip %s > %t.dev
 // RUN: cat %t.dev | FileCheck -check-prefixes=COMMON,DEV,RDC-D %s

 // DEV-DAG: @ex.managed = external addrspace(1) global i32, align 4
 // DEV-DAG: @ex = external addrspace(1) externally_initialized global i32 addrspace(1)*
 // HOST-DAG: @ex.managed = external global i32
 // HOST-DAG: @ex = external externally_initialized global i32*
 extern __managed__ int ex;
 
 // NORDC-D-DAG: @_ZL2sx.managed = addrspace(1) externally_initialized global i32 1, align 4
 // NORDC-D-DAG: @_ZL2sx = addrspace(1) externally_initialized global i32 addrspace(1)* null
-// RDC-D-DAG: @_ZL2sx__static__[[HASH:.*]].managed = addrspace(1) externally_initialized global i32 1, align 4
-// RDC-D-DAG: @_ZL2sx__static__[[HASH]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
+// RDC-D-DAG: @_ZL2sx.static.[[HASH:.*]].managed = addrspace(1) externally_initialized global i32 1, align 4
+// RDC-D-DAG: @_ZL2sx.static.[[HASH]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
 // HOST-DAG: @_ZL2sx.managed = internal global i32 1
 // HOST-DAG: @_ZL2sx = internal externally_initialized global i32* null
 // NORDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx\00"
-// RDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx__static__[[HASH:.*]]\00"
+// RDC-DAG: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH:.*]]\00"
 
-// POSTFIX:  @_ZL2sx__static__[[HASH:.*]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
-// POSTFIX: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx__static__[[HASH]]\00"
+// POSTFIX:  @_ZL2sx.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 addrspace(1)* null
+// POSTFIX: @[[DEVNAMESX:[0-9]+]] = {{.*}}c"_ZL2sx.static.[[HASH]]\00"
 static __managed__ int sx = 1;
 
 // DEV-DAG: @llvm.compiler.used
 // DEV-SAME-DAG: @x.managed
 // DEV-SAME-DAG: @x
 // DEV-SAME-DAG: @v.managed
 // DEV-SAME-DAG: @v
 // DEV-SAME-DAG: @_ZL2sx.managed

clang/test/CodeGenCUDA/static-device-var-rdc.cu

 // RUN: %clang_cc1 -no-opaque-pointers -triple amdgcn-amd-amdhsa -fcuda-is-device -cuid=abc \
 // RUN:   -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.dev
 // RUN: cat %t.dev | FileCheck -check-prefix=DEV-NEG %s
 
 // RUN: %clang_cc1 -no-opaque-pointers -triple x86_64-gnu-linux -cuid=abc \
 // RUN:   -std=c++11 -fgpu-rdc -emit-llvm -o - -x hip %s > %t.host
 // RUN: cat %t.host | FileCheck -check-prefix=HOST-NEG %s
 
+// Check postfix for CUDA.
+
+// RUN: %clang_cc1 -no-opaque-pointers -triple nvptx -fcuda-is-device -cuid=abc \
+// RUN:   -std=c++11 -fgpu-rdc -emit-llvm -o - %s | FileCheck \
+// RUN:   -check-prefixes=CUDA %s
 
 #include "Inputs/cuda.h"
 
 // Test function scope static device variable, which should not be externalized.
 // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
 
 
 // HOST-DAG: @_ZL1x = internal global i32 undef
 // HOST-DAG: @_ZL1y = internal global i32 undef
 
 // Test normal static device variables
 // INT-DEV-DAG: @_ZL1x = addrspace(1) externally_initialized global i32 0
 // INT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x\00"
 
 // Test externalized static device variables
-// EXT-DEV-DAG: @_ZL1x__static__[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
-// EXT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x__static__[[HASH:.*]]\00"
+// EXT-DEV-DAG: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
+// EXT-HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH:.*]]\00"
+// CUDA-DAG: @_ZL1x__static__[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
 
-// POSTFIX: @_ZL1x__static__[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
-// POSTFIX: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x__static__[[HASH]]\00"
+// POSTFIX: @_ZL1x.static.[[HASH:.*]] = addrspace(1) externally_initialized global i32 0
+// POSTFIX: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x.static.[[HASH]]\00"
 
 static __device__ int x;
 
 // Test static device variables not used by host code should not be externalized
 // DEV-DAG: @_ZL2x2 = internal addrspace(1) global i32 0
 
 static __device__ int x2;
 
 // Test normal static device variables
 // INT-DEV-DAG: @_ZL1y = addrspace(4) externally_initialized global i32 0
 // INT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y\00"
 
 // Test externalized static device variables
-// EXT-DEV-DAG: @_ZL1y__static__[[HASH]] = addrspace(4) externally_initialized global i32 0
-// EXT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y__static__[[HASH]]\00"
+// EXT-DEV-DAG: @_ZL1y.static.[[HASH]] = addrspace(4) externally_initialized global i32 0
+// EXT-HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y.static.[[HASH]]\00"
 
 static __constant__ int y;
 
 // Test static host variable, which should not be externalized nor registered.
 // HOST-DAG: @_ZL1z = internal global i32 0
 // DEV-NEG-NOT: @_ZL1z
 static int z;