[CUDA] Make sure we emit all templated __global__ functions on device side.

[tra]

Summary

Templated kernels that were instantiated from the host code would normally be eliminated because they were never referenced on device side.
Add implicit 'used' attribute to global functions which prevents their elimination.

[majnemer]

Couldn't you just add an implicit UsedAttr when processing the CUDAGlobalAttr
and LangOpts.CUDAIsDevice was set to true?

[tra]

Changed to use implicit 'used' attribute.

[tra]

Ping.

[majnemer]

LGTM

cfe/trunk/lib/Sema/SemaDeclAttr.cpp

S.Diag(FD->getTypeSpecStartLoc(), diag::err_kern_type_not_void_return)
         << (RTRange.isValid() ? FixItHint::CreateReplacement(RTRange, "void")
                               : 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;
   }

cfe/trunk/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__ functiona 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_
+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}