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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits