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

[tra]

Summary

This is a somewhat different way to do it than D11666 which got rolled back.

Codegen postpones emitting instantiated kernel function template until it's used.
If kernel is used only from the host side (which is normally the case) we'll never emit
it because on device side we don't emit the host code that uses it.

The change allows CUDA kernels to be emitted on device side unconditionally.
It's overly conservative and may emit more functions than we really need, but it
guarantees that the kernels launched from the host side are do exist on device-side.
In case it ever causes issues, there are other ways to address the issue,
though they are more invasive and are currently not worth the trouble.

[echristo]

LGTM. Thanks for working on this.

-eric

[eliben]

lgtm

test/CodeGenCUDA/ptx-kernels.cu

26

It't not clear what this metadata is part of? I'm guessing llvm.used, so maybe make that explicit with an earlier CHECK?

[tra]

Emitting IR is not sufficient to ensure that the kernels survive GDCE, so the patch does not work with optimizations on.
D11666 would have to do for now.

lib/AST/ASTContext.cpp

bool ASTContext::DeclMustBeEmitted(const Decl *D) {
   if (const FunctionDecl *FD = dyn_cast<FunctionDecl>(D)) {
     // Forward declarations aren't required.
     if (!FD->doesThisDeclarationHaveABody())
       return FD->doesDeclarationForceExternallyVisibleDefinition();
 
     // Constructors and destructors are required.
     if (FD->hasAttr<ConstructorAttr>() || FD->hasAttr<DestructorAttr>())
       return true;
-    
+    // Force all CUDA kernels to be emitted on device side.
+    // Otherwise, templated kernels may never be emitted as they are
+    // only used from host-side code which we never emit on device
+    // side and which therefore would never trigger us to emit
+    // device-side kernel it might've instantiated. The trade-off is
+    // that emitting all kernels is over-conservative and we may emit
+    // more of them than necessary. If excess of generated GPU code
+    // becomes a problem we can revisit this.
+    if (getLangOpts().CUDA && getLangOpts().CUDAIsDevice &&
+        FD->hasAttr<CUDAGlobalAttr>())
+      return true;
+
     // The key function for a class is required.  This rule only comes
     // into play when inline functions can be key functions, though.
     if (getTargetInfo().getCXXABI().canKeyFunctionBeInline()) {
       if (const CXXMethodDecl *MD = dyn_cast<CXXMethodDecl>(FD)) {
         const CXXRecordDecl *RD = MD->getParent();
         if (MD->isOutOfLine() && RD->isDynamicClass()) {
           const CXXMethodDecl *KeyFunc = getCurrentKeyFunction(RD);
           if (KeyFunc && KeyFunc->getCanonicalDecl() == MD->getCanonicalDecl())

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"
 
 // 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}