[CUDA] __global__ functions should always be visible externally.

[tra]

Summary

Adjusted global functions with DiscardableODR linkage to use StrongODR linkage instead so they are visible externally.

Replaces D11666 / r248293

[echristo]

Seems to be what we were chatting about.

Thanks!

-eric

lib/AST/ASTContext.cpp

static GVALinkage basicGVALinkageForFunction(const ASTContext &Context,
   // 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 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;
   }
   return L;
 }
 
 GVALinkage ASTContext::GetGVALinkageForFunction(const FunctionDecl *FD) const {
-  return adjustGVALinkageForDLLAttribute(basicGVALinkageForFunction(*this, FD),
-                                         FD);
+  return adjustGVALinkageForAttributes(basicGVALinkageForFunction(*this, FD),
 }
 
 static GVALinkage basicGVALinkageForVariable(const ASTContext &Context,
                                              const VarDecl *VD) {
   if (!VD->isExternallyVisible())
     return GVA_Internal;
 
   if (VD->isStaticLocal()) {
static GVALinkage basicGVALinkageForVariable(const ASTContext &Context,
   case TSK_ImplicitInstantiation:
     return GVA_DiscardableODR;
   }
 
   llvm_unreachable("Invalid Linkage!");
 }
 
 GVALinkage ASTContext::GetGVALinkageForVariable(const VarDecl *VD) {
-  return adjustGVALinkageForDLLAttribute(basicGVALinkageForVariable(*this, VD),
-                                         VD);
+  return adjustGVALinkageForAttributes(basicGVALinkageForVariable(*this, VD),
 }
 
 bool ASTContext::DeclMustBeEmitted(const Decl *D) {
   if (const VarDecl *VD = dyn_cast<VarDecl>(D)) {
     if (!VD->isFileVarDecl())
       return false;
     // Global named register variables (GNU extension) are never emitted.
     if (VD->getStorageClass() == SC_Register)

lib/Sema/SemaDeclAttr.cpp

S.Diag(FD->getTypeSpecStartLoc(), diag::err_kern_type_not_void_return)
                               : FixItHint());
     return;
   }
 
   D->addAttr(::new (S.Context)
               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) {
   FunctionDecl *Fn = cast<FunctionDecl>(D);
   if (!Fn->isInlineSpecified()) {
     S.Diag(Attr.getLoc(), diag::warn_gnu_inline_attribute_requires_inline);
     return;
   }

test/CodeGenCUDA/ptx-kernels.cu

 // Make sure that __global__ functions are emitted along with correct
 // annotations and are added to @llvm.used to prevent their elimination.
 // REQUIRES: nvptx-registered-target
 //
 // RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
 
 #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() {}
 
 // CHECK-LABEL: define void @global_function
 extern "C"
 __global__ void global_function() {
   // CHECK: call void @device_function
   device_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}
 // CHECK: !{{[0-9]+}} = !{void (i32)* @_Z16templated_kernelIiEvT_, !"kernel", i32 1}