r248400 - [CUDA] __global__ functions should always be visible externally.

[Artem Belevich via cfe-commits]

Author: tra
Date: Wed Sep 23 12:44:53 2015
New Revision: 248400

URL: http://llvm.org/viewvc/llvm-project?rev=248400&view=rev
Log:
[CUDA] __global__ functions should always be visible externally.

Adjust __global__ functions with DiscardableODR linkage to use
StrongODR linkage instead, so they are visible externally.

Differential Revision: http://reviews.llvm.org/D13067

Modified:
    cfe/trunk/lib/AST/ASTContext.cpp
    cfe/trunk/lib/Sema/SemaDeclAttr.cpp
    cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu

Modified: cfe/trunk/lib/AST/ASTContext.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTContext.cpp?rev=248400&r1=248399&r2=248400&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ASTContext.cpp (original)
+++ cfe/trunk/lib/AST/ASTContext.cpp Wed Sep 23 12:44:53 2015
@@ -8284,13 +8284,13 @@ static GVALinkage basicGVALinkageForFunc
   return GVA_DiscardableODR;
 }
 
-static GVALinkage adjustGVALinkageForDLLAttribute(GVALinkage L, const Decl *D) {
+static GVALinkage adjustGVALinkageForAttributes(GVALinkage L, const Decl *D) {
   // 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>()) {
+  } else if (D->hasAttr<DLLExportAttr>() || D->hasAttr<CUDAGlobalAttr>()) {
     if (L == GVA_DiscardableODR)
       return GVA_StrongODR;
   }
@@ -8298,8 +8298,8 @@ static GVALinkage adjustGVALinkageForDLL
 }
 
 GVALinkage ASTContext::GetGVALinkageForFunction(const FunctionDecl *FD) const {
-  return adjustGVALinkageForDLLAttribute(basicGVALinkageForFunction(*this, FD),
-                                         FD);
+  return adjustGVALinkageForAttributes(basicGVALinkageForFunction(*this, FD),
+                                       FD);
 }
 
 static GVALinkage basicGVALinkageForVariable(const ASTContext &Context,
@@ -8355,8 +8355,8 @@ static GVALinkage basicGVALinkageForVari
 }
 
 GVALinkage ASTContext::GetGVALinkageForVariable(const VarDecl *VD) {
-  return adjustGVALinkageForDLLAttribute(basicGVALinkageForVariable(*this, VD),
-                                         VD);
+  return adjustGVALinkageForAttributes(basicGVALinkageForVariable(*this, VD),
+                                       VD);
 }
 
 bool ASTContext::DeclMustBeEmitted(const Decl *D) {

Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=248400&r1=248399&r2=248400&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Wed Sep 23 12:44:53 2015
@@ -3351,9 +3351,6 @@ static void handleGlobalAttr(Sema &S, De
               CUDAGlobalAttr(Attr.getRange(), S.Context,
                              Attr.getAttributeSpellingListIndex()));
 
-  // Add implicit attribute((used)) so we don't eliminate kernels
-  // because there is nothing referencing them on device side.
-  D->addAttr(UsedAttr::CreateImplicit(S.Context));
 }
 
 static void handleGNUInlineAttr(Sema &S, Decl *D, const AttributeList &Attr) {

Modified: cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu?rev=248400&r1=248399&r2=248400&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu (original)
+++ cfe/trunk/test/CodeGenCUDA/ptx-kernels.cu Wed Sep 23 12:44:53 2015
@@ -6,11 +6,6 @@
 
 #include "Inputs/cuda.h"
 
-// Make sure that all __global__ functions are added to @llvm.used
-// CHECK: @llvm.used = appending global
-// CHECK-SAME: @global_function
-// CHECK-SAME: @_Z16templated_kernelIiEvT_
-
 // CHECK-LABEL: define void @device_function
 extern "C"
 __device__ void device_function() {}
@@ -24,7 +19,7 @@ __global__ void global_function() {
 
 // Make sure host-instantiated kernels are preserved on device side.
 template <typename T> __global__ void templated_kernel(T param) {}
-// CHECK-LABEL: define linkonce_odr void @_Z16templated_kernelIiEvT_
+// CHECK-LABEL: define weak_odr void @_Z16templated_kernelIiEvT_
 void host_function() { templated_kernel<<<0,0>>>(0); }
 
 // CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}