[OPENMP] Codegen for `omp declare target` construct.

Added initial codegen for device side of declarations inside `omp
declare target` construct + codegen for implicit `declare target`
functions, which are used in the target regions.

llvm-svn: 327636

clang/lib/AST/ASTContext.cpp

@@ -9402,8 +9402,7 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
       return false;
   } else if (isa<PragmaCommentDecl>(D))
     return true;
-  else if (isa<OMPThreadPrivateDecl>(D) ||
-           D->hasAttr<OMPDeclareTargetDeclAttr>())
+  else if (isa<OMPThreadPrivateDecl>(D))
     return true;
   else if (isa<PragmaDetectMismatchDecl>(D))
     return true;
@@ -9492,6 +9491,12 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
         if (DeclMustBeEmitted(BindingVD))
           return true;
 
+  // If the decl is marked as `declare target`, it should be emitted.
+  for (const auto *Decl = D->getMostRecentDecl(); Decl;
+       Decl = Decl->getPreviousDecl())
+    if (Decl->hasAttr<OMPDeclareTargetDeclAttr>())
+      return true;
+
   return false;
 }
 

clang/lib/CodeGen/CGDecl.cpp

@@ -285,8 +285,11 @@ llvm::Constant *CodeGenModule::getOrCreateStaticVarDecl(
     // never defer them.
     assert(isa<ObjCMethodDecl>(DC) && "unexpected parent code decl");
   }
-  if (GD.getDecl())
+  if (GD.getDecl()) {
+    // Disable emission of the parent function for the OpenMP device codegen.
+    CGOpenMPRuntime::DisableAutoDeclareTargetRAII NoDeclTarget(*this);
     (void)GetAddrOfGlobal(GD);
+  }
 
   return Addr;
 }

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -7405,9 +7405,14 @@ bool CGOpenMPRuntime::emitTargetFunctions(GlobalDecl GD) {
   // Try to detect target regions in the function.
   scanForTargetRegionsFunctions(FD.getBody(), CGM.getMangledName(GD));
 
-  // We should not emit any function other that the ones created during the
-  // scanning. Therefore, we signal that this function is completely dealt
-  // with.
+  // Do not to emit function if it is not marked as declare target.
+  if (!GD.getDecl()->hasAttrs())
+    return true;
+
+  for (const auto *D = FD.getMostRecentDecl(); D; D = D->getPreviousDecl())
+    if (D->hasAttr<OMPDeclareTargetDeclAttr>())
+      return false;
+
   return true;
 }
 
@@ -7433,8 +7438,15 @@ bool CGOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl GD) {
     }
   }
 
-  // If we are in target mode, we do not emit any global (declare target is not
-  // implemented yet). Therefore we signal that GD was processed in this case.
+  // Do not to emit variable if it is not marked as declare target.
+  if (!GD.getDecl()->hasAttrs())
+    return true;
+
+  for (const Decl *D = GD.getDecl()->getMostRecentDecl(); D;
+       D = D->getPreviousDecl())
+    if (D->hasAttr<OMPDeclareTargetDeclAttr>())
+      return false;
+
   return true;
 }
 
@@ -7446,6 +7458,38 @@ bool CGOpenMPRuntime::emitTargetGlobal(GlobalDecl GD) {
   return emitTargetGlobalVariable(GD);
 }
 
+CGOpenMPRuntime::DisableAutoDeclareTargetRAII::DisableAutoDeclareTargetRAII(
+    CodeGenModule &CGM)
+    : CGM(CGM) {
+  if (CGM.getLangOpts().OpenMPIsDevice) {
+    SavedShouldMarkAsGlobal = CGM.getOpenMPRuntime().ShouldMarkAsGlobal;
+    CGM.getOpenMPRuntime().ShouldMarkAsGlobal = false;
+  }
+}
+
+CGOpenMPRuntime::DisableAutoDeclareTargetRAII::~DisableAutoDeclareTargetRAII() {
+  if (CGM.getLangOpts().OpenMPIsDevice)
+    CGM.getOpenMPRuntime().ShouldMarkAsGlobal = SavedShouldMarkAsGlobal;
+}
+
+bool CGOpenMPRuntime::markAsGlobalTarget(const FunctionDecl *D) {
+  if (!CGM.getLangOpts().OpenMPIsDevice || !ShouldMarkAsGlobal)
+    return true;
+  // Do not to emit function if it is marked as declare target as it was already
+  // emitted.
+  for (const auto *FD = D->getMostRecentDecl(); FD; FD = FD->getPreviousDecl())
+    if (FD->hasAttr<OMPDeclareTargetDeclAttr>())
+      return true;
+
+  const FunctionDecl *FD = D->getCanonicalDecl();
+  // Do not mark member functions except for static.
+  if (const auto *Method = dyn_cast<CXXMethodDecl>(FD))
+    if (!Method->isStatic())
+      return true;
+
+  return !AlreadyEmittedTargetFunctions.insert(FD).second;
+}
+
 llvm::Function *CGOpenMPRuntime::emitRegistrationFunction() {
   // If we have offloading in the current module, we need to emit the entries
   // now and register the offloading descriptor.

clang/lib/CodeGen/CGOpenMPRuntime.h

@@ -199,6 +199,18 @@ public:
 };
 
 class CGOpenMPRuntime {
+public:
+  /// Allows to disable automatic handling of functions used in target regions
+  /// as those marked as `omp declare target`.
+  class DisableAutoDeclareTargetRAII {
+    CodeGenModule &CGM;
+    bool SavedShouldMarkAsGlobal;
+
+  public:
+    DisableAutoDeclareTargetRAII(CodeGenModule &CGM);
+    ~DisableAutoDeclareTargetRAII();
+  };
+
 protected:
   CodeGenModule &CGM;
 
@@ -488,6 +500,9 @@ private:
   };
   OffloadEntriesInfoManagerTy OffloadEntriesInfoManager;
 
+  bool ShouldMarkAsGlobal = true;
+  llvm::SmallDenseSet<const FunctionDecl *> AlreadyEmittedTargetFunctions;
+
   /// \brief Creates and registers offloading binary descriptor for the current
   /// compilation unit. The function that does the registration is returned.
   llvm::Function *createOffloadingBinaryDescriptorRegistration();
@@ -1370,6 +1385,11 @@ public:
   /// Gets the OpenMP-specific address of the local variable.
   virtual Address getAddressOfLocalVariable(CodeGenFunction &CGF,
                                             const VarDecl *VD);
+
+  /// Marks the declaration as alread emitted for the device code and returns
+  /// true, if it was marked already, and false, otherwise.
+  bool markAsGlobalTarget(const FunctionDecl *D);
+
 };
 
 /// Class supports emissionof SIMD-only code.

clang/lib/CodeGen/CGStmtOpenMP.cpp

@@ -3914,6 +3914,16 @@ static void emitCommonOMPTargetDirective(CodeGenFunction &CGF,
   assert(isOpenMPTargetExecutionDirective(S.getDirectiveKind()));
   CodeGenModule &CGM = CGF.CGM;
 
+  // On device emit this construct as inlined code.
+  if (CGM.getLangOpts().OpenMPIsDevice) {
+    OMPLexicalScope Scope(CGF, S, OMPD_target);
+    CGM.getOpenMPRuntime().emitInlinedDirective(
+        CGF, OMPD_target, [&S](CodeGenFunction &CGF, PrePostActionTy &) {
+          CGF.EmitStmt(S.getCapturedStmt(OMPD_target)->getCapturedStmt());
+        });
+    return;
+  }
+
   llvm::Function *Fn = nullptr;
   llvm::Constant *FnID = nullptr;
 

clang/lib/CodeGen/CodeGenModule.cpp

@@ -2383,6 +2383,12 @@ llvm::Constant *CodeGenModule::GetOrCreateLLVMFunction(
   // Any attempts to use a MultiVersion function should result in retrieving
   // the iFunc instead. Name Mangling will handle the rest of the changes.
   if (const FunctionDecl *FD = cast_or_null<FunctionDecl>(D)) {
+    // For the device mark the function as one that should be emitted.
+    if (getLangOpts().OpenMPIsDevice && OpenMPRuntime &&
+        !OpenMPRuntime->markAsGlobalTarget(FD) && FD->isDefined() &&
+        !DontDefer && !IsForDefinition)
+      addDeferredDeclToEmit(GD);
+
     if (FD->isMultiVersion() && FD->getAttr<TargetAttr>()->isDefaultVersion()) {
       UpdateMultiVersionNames(GD, FD);
       if (!IsForDefinition)
@@ -3072,6 +3078,12 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   if (getLangOpts().OpenCL && ASTTy->isSamplerT())
     return;
 
+  // If this is OpenMP device, check if it is legal to emit this global
+  // normally.
+  if (LangOpts.OpenMPIsDevice && OpenMPRuntime &&
+      OpenMPRuntime->emitTargetGlobalVariable(D))
+    return;
+
   llvm::Constant *Init = nullptr;
   CXXRecordDecl *RD = ASTTy->getBaseElementTypeUnsafe()->getAsCXXRecordDecl();
   bool NeedsGlobalCtor = false;

clang/lib/Parse/ParseOpenMP.cpp

@@ -758,6 +758,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
     if (!Actions.ActOnStartOpenMPDeclareTargetDirective(DTLoc))
       return DeclGroupPtrTy();
 
+    llvm::SmallVector<Decl *, 4>  Decls;
     DKind = ParseOpenMPDirectiveKind(*this);
     while (DKind != OMPD_end_declare_target && DKind != OMPD_declare_target &&
            Tok.isNot(tok::eof) && Tok.isNot(tok::r_brace)) {
@@ -772,6 +773,10 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
         Ptr =
             ParseCXXClassMemberDeclarationWithPragmas(AS, Attrs, TagType, Tag);
       }
+      if (Ptr) {
+        DeclGroupRef Ref = Ptr.get();
+        Decls.append(Ref.begin(), Ref.end());
+      }
       if (Tok.isAnnotation() && Tok.is(tok::annot_pragma_openmp)) {
         TentativeParsingAction TPA(*this);
         ConsumeAnnotationToken();
@@ -797,7 +802,8 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
       Diag(DTLoc, diag::note_matching) << "'#pragma omp declare target'";
     }
     Actions.ActOnFinishOpenMPDeclareTargetDirective();
-    return DeclGroupPtrTy();
+    return DeclGroupPtrTy::make(DeclGroupRef::Create(
+        Actions.getASTContext(), Decls.begin(), Decls.size()));
   }
   case OMPD_unknown:
     Diag(Tok, diag::err_omp_unknown_directive);

clang/lib/Sema/SemaOpenMP.cpp

@@ -1382,13 +1382,17 @@ VarDecl *Sema::IsOpenMPCapturedDecl(ValueDecl *D) {
   // If we are attempting to capture a global variable in a directive with
   // 'target' we return true so that this global is also mapped to the device.
   //
-  // FIXME: If the declaration is enclosed in a 'declare target' directive,
-  // then it should not be captured. Therefore, an extra check has to be
-  // inserted here once support for 'declare target' is added.
-  //
   auto *VD = dyn_cast<VarDecl>(D);
-  if (VD && !VD->hasLocalStorage() && isInOpenMPTargetExecutionDirective())
+  if (VD && !VD->hasLocalStorage() && isInOpenMPTargetExecutionDirective()) {
+    // If the declaration is enclosed in a 'declare target' directive,
+    // then it should not be captured.
+    //
+    for (const auto *Var = VD->getMostRecentDecl(); Var;
+         Var = Var->getPreviousDecl())
+      if (Var->hasAttr<OMPDeclareTargetDeclAttr>())
+        return nullptr;
     return VD;
+  }
 
   if (DSAStack->getCurrentDirective() != OMPD_unknown &&
       (!DSAStack->isClauseParsingMode() ||

clang/test/OpenMP/declare_target_codegen.cpp

@@ -0,0 +1,66 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o -| FileCheck %s --check-prefix SIMD-ONLY
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -o - | FileCheck %s --check-prefix SIMD-ONLY
+
+// expected-no-diagnostics
+
+// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
+
+// CHECK-NOT: define {{.*}}{{baz1|baz4|maini1}}
+// CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,
+// CHECK-DAG: @b = global i32 15,
+// CHECK-DAG: @d = global i32 0,
+// CHECK-DAG: @c = external global i32,
+
+// CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3}}{{.*}}()
+
+#ifndef HEADER
+#define HEADER
+
+int foo();
+
+int baz1();
+
+int baz2();
+
+int baz4() { return 5; }
+
+#pragma omp declare target
+int foo() { return 0; }
+int b = 15;
+int d;
+#pragma omp end declare target
+int c;
+
+int bar() { return 1 + foo() + bar() + baz1() + baz2(); }
+
+int maini1() {
+  int a;
+  static long aa = 32;
+// CHECK-DAG: define void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* dereferenceable{{.*}}, i64 {{.*}}, i64 {{.*}})
+#pragma omp target map(tofrom \
+                       : a)
+  {
+    static long aaa = 23;
+    a = foo() + bar() + b + c + d + aa + aaa;
+  }
+  return baz4();
+}
+
+int baz3();
+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