[OPENMP] Emit template instatiation|specialization functions for

devices.

If the function is an instantiation|specialization of the template and
is used in the device code, the definitions of such functions should be
emitted for the device.

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@331261 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp
index 0500f23b3d9b49edd9c2abc19e5da18099941c54..b7e4cb46c89eba08e6c51690e67fa279269bdb77 100644 (file)
--- a/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -907,6 +907,14 @@ isDeclareTargetDeclaration(const ValueDecl *VD) {
     if (const auto *Attr = D->getAttr<OMPDeclareTargetDeclAttr>())
       return Attr->getMapType();
   }
+  if (const auto *V = dyn_cast<VarDecl>(VD)) {
+    if (const VarDecl *TD = V->getTemplateInstantiationPattern())
+      return isDeclareTargetDeclaration(TD);
+  } else if (const auto *FD = dyn_cast<FunctionDecl>(VD)) {
+    if (const auto *TD = FD->getTemplateInstantiationPattern())
+      return isDeclareTargetDeclaration(TD);
+  }
+
   return llvm::None;
 }
 
@@ -7795,7 +7803,8 @@ bool CGOpenMPRuntime::emitTargetFunctions(GlobalDecl GD) {
   scanForTargetRegionsFunctions(FD->getBody(), CGM.getMangledName(GD));
 
   // Do not to emit function if it is not marked as declare target.
-  return !isDeclareTargetDeclaration(FD);
+  return !isDeclareTargetDeclaration(FD) &&
+         AlreadyEmittedTargetFunctions.count(FD->getCanonicalDecl()) == 0;
 }
 
 bool CGOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl GD) {

diff --git a/test/OpenMP/declare_target_codegen.cpp b/test/OpenMP/declare_target_codegen.cpp
index a5a51bc26f0d95d94cf5da88090a6ba5cae0a362..15b9f58833d332e8ccca1cc471d372582d82759d 100644 (file)
--- a/test/OpenMP/declare_target_codegen.cpp
+++ b/test/OpenMP/declare_target_codegen.cpp
@@ -18,7 +18,7 @@
 // CHECK-DAG: @d = global i32 0,
 // CHECK-DAG: @c = external global i32,
 
-// CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3}}{{.*}}()
+// CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA}}{{.*}}()
 
 #ifndef HEADER
 #define HEADER
@@ -31,6 +31,11 @@ int baz2();
 
 int baz4() { return 5; }
 
+template <typename T>
+T FA() {
+  return T();
+}
+
 #pragma omp declare target
 struct S {
   int a;
@@ -54,19 +59,18 @@ int maini1() {
   {
     S s(a);
     static long aaa = 23;
-    a = foo() + bar() + b + c + d + aa + aaa;
+    a = foo() + bar() + b + c + d + aa + aaa + FA<int>();
   }
   return baz4();
 }
 
-int baz3();
+int baz3() { return 2 + baz2(); }
 int baz2() {
 // CHECK-DAG: define void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}})
 #pragma omp target
   ++c;
   return 2 + baz3();
 }
-int baz3() { return 2 + baz2(); }
 
 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1}}
 #endif // HEADER