From 09795e30e6ae4fbb8663b50738d9bd66f097f128 Mon Sep 17 00:00:00 2001
From: Jonas Hahnfeld <Hahnfeld@itc.rwth-aachen.de>
Date: Wed, 15 Jul 2015 15:56:34 +0200
Subject: [PATCH] Initialize target library with only a target data

This corner case may occur in benchmarks that want to measure data movement.
---
 lib/CodeGen/CGOpenMPRuntime.cpp |  4 ++--
 lib/CodeGen/CGOpenMPRuntime.h   | 11 ++++++++++-
 lib/CodeGen/CGStmtOpenMP.cpp    |  1 +
 3 files changed, 13 insertions(+), 3 deletions(-)

diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp
index 32fe43ce15..06f2d4cf44 100644
--- a/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -81,8 +81,8 @@ void CGOpenMPRuntime::registerEntryForDeclaration(const Decl *D,
 
 CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM)
     : CGM(CGM), DefaultOpenMPPSource(nullptr), NumTargetRegions(0),
-      NumTargetGlobals(0), HasTargetInfoLoaded(false),
-      TargetRegionsDescriptor(nullptr) {
+      NumTargetGlobals(0), HasTargetDataRegions(false),
+      HasTargetInfoLoaded(false), TargetRegionsDescriptor(nullptr) {
   IdentTy = llvm::StructType::create(
       "ident_t", CGM.Int32Ty /* reserved_1 */, CGM.Int32Ty /* flags */,
       CGM.Int32Ty /* reserved_2 */, CGM.Int32Ty /* reserved_3 */,
diff --git a/lib/CodeGen/CGOpenMPRuntime.h b/lib/CodeGen/CGOpenMPRuntime.h
index d4a41dd121..4f276f0829 100644
--- a/lib/CodeGen/CGOpenMPRuntime.h
+++ b/lib/CodeGen/CGOpenMPRuntime.h
@@ -163,6 +163,9 @@ class CGOpenMPRuntime {
   // Number of globals processed so far that are to be mapped into a target
   unsigned NumTargetGlobals;
 
+  // True if target data region was processed
+  bool HasTargetDataRegions;
+
   // Name of the current function whose target regions are being identified
   std::string CurTargetParentFunctionName;
 
@@ -284,10 +287,16 @@ class CGOpenMPRuntime {
   // entry point
   bool isValidOtherTargetFunction(StringRef name);
 
+  // Register target data region
+  void setHasTargetDataRegions(bool val) {
+    HasTargetDataRegions = val;
+  }
+
   // Return true if the current module requires a the target descriptor to be
   // registered
   bool requiresTargetDescriptorRegistry(){
-    return NumTargetRegions != 0 || !TargetGlobalInitializers.empty();
+    return NumTargetRegions != 0 || !TargetGlobalInitializers.empty() ||
+           HasTargetDataRegions;
   }
 
   // Register global initializer for OpenMP Target offloading
diff --git a/lib/CodeGen/CGStmtOpenMP.cpp b/lib/CodeGen/CGStmtOpenMP.cpp
index d97c81cda6..e85bbb85c5 100644
--- a/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/lib/CodeGen/CGStmtOpenMP.cpp
@@ -6864,6 +6864,7 @@ void CodeGenFunction::EmitOMPTargetDataDirective(
       EmitFinalOMPClause(*(*I), S);
 
   CGM.OpenMPSupport.endOpenMPRegion();
+  CGM.getOpenMPRuntime().setHasTargetDataRegions(true);
 }
 
 // Generate the instructions for '#pragma omp target update' directive.