diff options
| -rw-r--r-- | clang/lib/Sema/SemaOpenMP.cpp | 48 | ||||
| -rw-r--r-- | clang/test/OpenMP/declare_target_link_codegen.cpp | 10 |
2 files changed, 53 insertions, 5 deletions
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 6c4c4a845ce..32e493ed57c 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -147,6 +147,9 @@ private: /// Reference to the taskgroup task_reduction reference expression. Expr *TaskgroupReductionRef = nullptr; llvm::DenseSet<QualType> MappedClassesQualTypes; + /// List of globals marked as declare target link in this target region + /// (isOpenMPTargetExecutionDirective(Directive) == true). + llvm::SmallVector<DeclRefExpr *, 4> DeclareTargetLinkVarDecls; SharingMapTy(OpenMPDirectiveKind DKind, DeclarationNameInfo Name, Scope *CurScope, SourceLocation Loc) : Directive(DKind), DirectiveName(Name), CurScope(CurScope), @@ -674,6 +677,31 @@ public: return StackElem.MappedClassesQualTypes.count(QT) != 0; } + /// Adds global declare target to the parent target region. + void addToParentTargetRegionLinkGlobals(DeclRefExpr *E) { + assert(*OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration( + E->getDecl()) == OMPDeclareTargetDeclAttr::MT_Link && + "Expected declare target link global."); + if (isStackEmpty()) + return; + auto It = Stack.back().first.rbegin(); + while (It != Stack.back().first.rend() && + !isOpenMPTargetExecutionDirective(It->Directive)) + ++It; + if (It != Stack.back().first.rend()) { + assert(isOpenMPTargetExecutionDirective(It->Directive) && + "Expected target executable directive."); + It->DeclareTargetLinkVarDecls.push_back(E); + } + } + + /// Returns the list of globals with declare target link if current directive + /// is target. + ArrayRef<DeclRefExpr *> getLinkGlobals() const { + assert(isOpenMPTargetExecutionDirective(getCurrentDirective()) && + "Expected target executable directive."); + return Stack.back().first.back().DeclareTargetLinkVarDecls; + } }; bool isImplicitTaskingRegion(OpenMPDirectiveKind DKind) { @@ -2414,8 +2442,18 @@ public: // Define implicit data-sharing attributes for task. DVar = Stack->getImplicitDSA(VD, /*FromParent=*/false); if (isOpenMPTaskingDirective(DKind) && DVar.CKind != OMPC_shared && - !Stack->isLoopControlVariable(VD).first) + !Stack->isLoopControlVariable(VD).first) { ImplicitFirstprivate.push_back(E); + return; + } + + // Store implicitly used globals with declare target link for parent + // target. + if (!isOpenMPTargetExecutionDirective(DKind) && Res && + *Res == OMPDeclareTargetDeclAttr::MT_Link) { + Stack->addToParentTargetRegionLinkGlobals(E); + return; + } } } void VisitMemberExpr(MemberExpr *E) { @@ -2573,7 +2611,13 @@ public: } DSAAttrChecker(DSAStackTy *S, Sema &SemaRef, CapturedStmt *CS) - : Stack(S), SemaRef(SemaRef), ErrorFound(false), CS(CS) {} + : Stack(S), SemaRef(SemaRef), ErrorFound(false), CS(CS) { + // Process declare target link variables for the target directives. + if (isOpenMPTargetExecutionDirective(S->getCurrentDirective())) { + for (DeclRefExpr *E : Stack->getLinkGlobals()) + Visit(E); + } + } }; } // namespace diff --git a/clang/test/OpenMP/declare_target_link_codegen.cpp b/clang/test/OpenMP/declare_target_link_codegen.cpp index 564581c2592..36c295b2de5 100644 --- a/clang/test/OpenMP/declare_target_link_codegen.cpp +++ b/clang/test/OpenMP/declare_target_link_codegen.cpp @@ -37,10 +37,13 @@ int maini1() { { a = c; } +#pragma omp target +#pragma omp teams + c = a; return 0; } -// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](i32* dereferenceable{{[^,]*}} +// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-10]](i32* dereferenceable{{[^,]*}} // DEVICE: [[C_REF:%.+]] = load i32*, i32** @c_decl_tgt_link_ptr, // DEVICE: [[C:%.+]] = load i32, i32* [[C_REF]], // DEVICE: store i32 [[C]], i32* % @@ -59,9 +62,10 @@ int maini1() { // HOST: [[BP0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // HOST: [[P0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // HOST: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 2, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0)) -// HOST: call void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-26]](i32* %{{[^,]+}}) +// HOST: call void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-29]](i32* %{{[^,]+}}) +// HOST: call i32 @__tgt_target_teams(i64 -1, i8* @.__omp_offloading_{{.+}}_l40.region_id, i32 2, {{.+}}) -// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-28]](i32* dereferenceable{{.*}}) +// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-32]](i32* dereferenceable{{.*}}) // HOST: [[C:%.*]] = load i32, i32* @c, // HOST: store i32 [[C]], i32* % |
