r327636 - [OPENMP] Codegen for `omp declare target` construct.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Thu Mar 15 08:47:20 PDT 2018
Author: abataev
Date: Thu Mar 15 08:47:20 2018
New Revision: 327636
URL: http://llvm.org/viewvc/llvm-project?rev=327636&view=rev
Log:
[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.
Added:
cfe/trunk/test/OpenMP/declare_target_codegen.cpp
Modified:
cfe/trunk/lib/AST/ASTContext.cpp
cfe/trunk/lib/CodeGen/CGDecl.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
cfe/trunk/lib/Parse/ParseOpenMP.cpp
cfe/trunk/lib/Sema/SemaOpenMP.cpp
Modified: cfe/trunk/lib/AST/ASTContext.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTContext.cpp?rev=327636&r1=327635&r2=327636&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ASTContext.cpp (original)
+++ cfe/trunk/lib/AST/ASTContext.cpp Thu Mar 15 08:47:20 2018
@@ -9402,8 +9402,7 @@ bool ASTContext::DeclMustBeEmitted(const
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
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;
}
Modified: cfe/trunk/lib/CodeGen/CGDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDecl.cpp?rev=327636&r1=327635&r2=327636&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDecl.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDecl.cpp Thu Mar 15 08:47:20 2018
@@ -285,8 +285,11 @@ llvm::Constant *CodeGenModule::getOrCrea
// 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;
}
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=327636&r1=327635&r2=327636&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Thu Mar 15 08:47:20 2018
@@ -7405,9 +7405,14 @@ bool CGOpenMPRuntime::emitTargetFunction
// 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::emitTargetGlobalVa
}
}
- // 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(G
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.
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h?rev=327636&r1=327635&r2=327636&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h Thu Mar 15 08:47:20 2018
@@ -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.
Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=327636&r1=327635&r2=327636&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Thu Mar 15 08:47:20 2018
@@ -3914,6 +3914,16 @@ static void emitCommonOMPTargetDirective
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;
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=327636&r1=327635&r2=327636&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Thu Mar 15 08:47:20 2018
@@ -2383,6 +2383,12 @@ llvm::Constant *CodeGenModule::GetOrCrea
// 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::EmitGlobalVarDefinit
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;
Modified: cfe/trunk/lib/Parse/ParseOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParseOpenMP.cpp?rev=327636&r1=327635&r2=327636&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/ParseOpenMP.cpp (original)
+++ cfe/trunk/lib/Parse/ParseOpenMP.cpp Thu Mar 15 08:47:20 2018
@@ -758,6 +758,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpen
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::ParseOpen
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::ParseOpen
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);
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=327636&r1=327635&r2=327636&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Thu Mar 15 08:47:20 2018
@@ -1382,13 +1382,17 @@ VarDecl *Sema::IsOpenMPCapturedDecl(Valu
// 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() ||
Added: cfe/trunk/test/OpenMP/declare_target_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_codegen.cpp?rev=327636&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/declare_target_codegen.cpp (added)
+++ cfe/trunk/test/OpenMP/declare_target_codegen.cpp Thu Mar 15 08:47:20 2018
@@ -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
More information about the cfe-commits
mailing list