From fb38828cb2834e7f685392d48ebcd5aa8a2ddc57 Mon Sep 17 00:00:00 2001
From: Alexey Bataev <a.bataev@hotmail.com>
Date: Tue, 1 May 2018 14:09:46 +0000
Subject: [PATCH] [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.

llvm-svn: 331261
---
 clang/lib/CodeGen/CGOpenMPRuntime.cpp        | 11 ++++++++++-
 clang/test/OpenMP/declare_target_codegen.cpp | 12 ++++++++----
 2 files changed, 18 insertions(+), 5 deletions(-)

diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
index 0500f23b3d9b..b7e4cb46c89e 100644
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/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/clang/test/OpenMP/declare_target_codegen.cpp b/clang/test/OpenMP/declare_target_codegen.cpp
index a5a51bc26f0d..15b9f58833d3 100644
--- a/clang/test/OpenMP/declare_target_codegen.cpp
+++ b/clang/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