diff options
| author | Patrick Lyster <Patrick.lyster@ibm.com> | 2019-01-02 19:28:48 +0000 |
|---|---|---|
| committer | Patrick Lyster <Patrick.lyster@ibm.com> | 2019-01-02 19:28:48 +0000 |
| commit | e13b1e3299c1feb642276491afc5a8cfec9e51b0 (patch) | |
| tree | c095757aa3d908dfe29d7dfe8830dedaf57ae57b | |
| parent | 4a401e9479a084e9037a47611f9004e0228c02f8 (diff) | |
| download | bcm5719-llvm-e13b1e3299c1feb642276491afc5a8cfec9e51b0.tar.gz bcm5719-llvm-e13b1e3299c1feb642276491afc5a8cfec9e51b0.zip | |
[OpenMP] Added support for explicit mapping of classes using 'this' pointer. Differential revision: https://reviews.llvm.org/D55982
llvm-svn: 350252
| -rw-r--r-- | clang/include/clang/Basic/DiagnosticSemaKinds.td | 8 | ||||
| -rw-r--r-- | clang/lib/CodeGen/CGOpenMPRuntime.cpp | 11 | ||||
| -rw-r--r-- | clang/lib/Sema/SemaOpenMP.cpp | 72 | ||||
| -rw-r--r-- | clang/test/OpenMP/target_ast_print.cpp | 35 | ||||
| -rw-r--r-- | clang/test/OpenMP/target_codegen.cpp | 50 | ||||
| -rw-r--r-- | clang/test/OpenMP/target_messages.cpp | 12 |
6 files changed, 183 insertions, 5 deletions
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 46cf0a423e4..c69566a2d22 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9076,6 +9076,14 @@ def note_omp_requires_previous_clause : Note < "%0 clause previously used here">; def err_omp_invalid_scope : Error < "'#pragma omp %0' directive must appear only in file scope">; +def note_omp_invalid_length_on_this_ptr_mapping : Note < + "expected length on mapping of 'this' array section expression to be '1'">; +def note_omp_invalid_lower_bound_on_this_ptr_mapping : Note < + "expected lower bound on mapping of 'this' array section expression to be '0' or not specified">; +def note_omp_invalid_subscript_on_this_ptr_map : Note < + "expected 'this' subscript expression on map clause to be 'this[0]'">; +def err_omp_invalid_map_this_expr : Error < + "invalid 'this' expression on 'map' clause">; } // end of OpenMP category let CategoryName = "Related Result Type Issue" in { diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 746d3d3e647..81760938406 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -6992,15 +6992,22 @@ private: // components. bool IsExpressionFirstInfo = true; Address BP = Address::invalid(); + const Expr *AssocExpr = I->getAssociatedExpression(); + const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr); + const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr); - if (isa<MemberExpr>(I->getAssociatedExpression())) { + if (isa<MemberExpr>(AssocExpr)) { // The base is the 'this' pointer. The content of the pointer is going // to be the base of the field being mapped. BP = CGF.LoadCXXThisAddress(); + } else if ((AE && isa<CXXThisExpr>(AE->getBase()->IgnoreParenImpCasts())) || + (OASE && + isa<CXXThisExpr>(OASE->getBase()->IgnoreParenImpCasts()))) { + BP = CGF.EmitOMPSharedLValue(AssocExpr).getAddress(); } else { // The base is the reference to the variable. // BP = &Var. - BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression()).getAddress(); + BP = CGF.EmitOMPSharedLValue(AssocExpr).getAddress(); if (const auto *VD = dyn_cast_or_null<VarDecl>(I->getAssociatedDeclaration())) { if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res = diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index b4eb4664763..78bef59ff6f 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -22,6 +22,7 @@ #include "clang/AST/StmtCXX.h" #include "clang/AST/StmtOpenMP.h" #include "clang/AST/StmtVisitor.h" +#include "clang/AST/TypeOrdering.h" #include "clang/Basic/OpenMPKinds.h" #include "clang/Sema/Initialization.h" #include "clang/Sema/Lookup.h" @@ -146,6 +147,7 @@ private: SourceLocation InnerTeamsRegionLoc; /// Reference to the taskgroup task_reduction reference expression. Expr *TaskgroupReductionRef = nullptr; + llvm::DenseSet<QualType> MappedClassesQualTypes; SharingMapTy(OpenMPDirectiveKind DKind, DeclarationNameInfo Name, Scope *CurScope, SourceLocation Loc) : Directive(DKind), DirectiveName(Name), CurScope(CurScope), @@ -660,6 +662,19 @@ public: return llvm::make_range(StackElem.DoacrossDepends.end(), StackElem.DoacrossDepends.end()); } + + // Store types of classes which have been explicitly mapped + void addMappedClassesQualTypes(QualType QT) { + SharingMapTy &StackElem = Stack.back().first.back(); + StackElem.MappedClassesQualTypes.insert(QT); + } + + // Return set of mapped classes types + bool isClassPreviouslyMapped(QualType QT) const { + const SharingMapTy &StackElem = Stack.back().first.back(); + return StackElem.MappedClassesQualTypes.count(QT) != 0; + } + }; bool isParallelOrTaskRegion(OpenMPDirectiveKind DKind) { return isOpenMPParallelDirective(DKind) || isOpenMPTaskingDirective(DKind) || @@ -2267,7 +2282,7 @@ public: return; auto *FD = dyn_cast<FieldDecl>(E->getMemberDecl()); OpenMPDirectiveKind DKind = Stack->getCurrentDirective(); - if (isa<CXXThisExpr>(E->getBase()->IgnoreParens())) { + if (auto *TE = dyn_cast<CXXThisExpr>(E->getBase()->IgnoreParens())) { if (!FD) return; DSAStackTy::DSAVarData DVar = Stack->getTopDSA(FD, /*FromParent=*/false); @@ -2294,6 +2309,12 @@ public: // if (FD->isBitField()) return; + + // Check to see if the member expression is referencing a class that + // has already been explicitly mapped + if (Stack->isClassPreviouslyMapped(TE->getType())) + return; + ImplicitMap.emplace_back(E); return; } @@ -12448,6 +12469,19 @@ static const Expr *checkMapClauseExpressionBase( E->getType())) AllowWholeSizeArraySection = false; + if (const auto *TE = dyn_cast<CXXThisExpr>(E)) { + Expr::EvalResult Result; + if (CurE->getIdx()->EvaluateAsInt(Result, SemaRef.getASTContext())) { + if (!Result.Val.getInt().isNullValue()) { + SemaRef.Diag(CurE->getIdx()->getExprLoc(), + diag::err_omp_invalid_map_this_expr); + SemaRef.Diag(CurE->getIdx()->getExprLoc(), + diag::note_omp_invalid_subscript_on_this_ptr_map); + } + } + RelevantExpr = TE; + } + // Record the component - we don't have any declaration associated. CurComponents.emplace_back(CurE, nullptr); } else if (auto *CurE = dyn_cast<OMPArraySectionExpr>(E)) { @@ -12494,6 +12528,30 @@ static const Expr *checkMapClauseExpressionBase( return nullptr; } + if (const auto *TE = dyn_cast<CXXThisExpr>(E)) { + Expr::EvalResult ResultR; + Expr::EvalResult ResultL; + if (CurE->getLength()->EvaluateAsInt(ResultR, + SemaRef.getASTContext())) { + if (!ResultR.Val.getInt().isOneValue()) { + SemaRef.Diag(CurE->getLength()->getExprLoc(), + diag::err_omp_invalid_map_this_expr); + SemaRef.Diag(CurE->getLength()->getExprLoc(), + diag::note_omp_invalid_length_on_this_ptr_mapping); + } + } + if (CurE->getLowerBound() && CurE->getLowerBound()->EvaluateAsInt( + ResultL, SemaRef.getASTContext())) { + if (!ResultL.Val.getInt().isNullValue()) { + SemaRef.Diag(CurE->getLowerBound()->getExprLoc(), + diag::err_omp_invalid_map_this_expr); + SemaRef.Diag(CurE->getLowerBound()->getExprLoc(), + diag::note_omp_invalid_lower_bound_on_this_ptr_mapping); + } + } + RelevantExpr = TE; + } + // Record the component - we don't have any declaration associated. CurComponents.emplace_back(CurE, nullptr); } else { @@ -12831,6 +12889,18 @@ checkMappableExpressionList(Sema &SemaRef, DSAStackTy *DSAS, assert(!CurComponents.empty() && "Invalid mappable expression information."); + if (const auto *TE = dyn_cast<CXXThisExpr>(BE)) { + // Add store "this" pointer to class in DSAStackTy for future checking + DSAS->addMappedClassesQualTypes(TE->getType()); + // Skip restriction checking for variable or field declarations + MVLI.ProcessedVarList.push_back(RE); + MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1); + MVLI.VarComponents.back().append(CurComponents.begin(), + CurComponents.end()); + MVLI.VarBaseDeclarations.push_back(nullptr); + continue; + } + // For the following checks, we rely on the base declaration which is // expected to be associated with the last component. The declaration is // expected to be a variable or a field (if 'this' is being mapped). diff --git a/clang/test/OpenMP/target_ast_print.cpp b/clang/test/OpenMP/target_ast_print.cpp index 2734294ff60..9b866b789a7 100644 --- a/clang/test/OpenMP/target_ast_print.cpp +++ b/clang/test/OpenMP/target_ast_print.cpp @@ -191,6 +191,41 @@ T tmain(T argc, T *argv) { // CHECK-NEXT: #pragma omp target defaultmap(tofrom: scalar) // CHECK-NEXT: foo() +// CHECK-LABEL: class S { +class S { + void foo() { +// CHECK-NEXT: void foo() { + int a = 0; +// CHECK-NEXT: int a = 0; + #pragma omp target map(this[0]) +// CHECK-NEXT: #pragma omp target map(tofrom: this[0]) + a++; +// CHECK-NEXT: a++; + #pragma omp target map(this[:1]) +// CHECK-NEXT: #pragma omp target map(tofrom: this[:1]) + a++; +// CHECK-NEXT: a++; + #pragma omp target map((this)[0]) +// CHECK-NEXT: #pragma omp target map(tofrom: (this)[0]) + a++; +// CHECK-NEXT: a++; + #pragma omp target map(this[:a]) +// CHECK-NEXT: #pragma omp target map(tofrom: this[:a]) + a++; +// CHECK-NEXT: a++; + #pragma omp target map(this[a:1]) +// CHECK-NEXT: #pragma omp target map(tofrom: this[a:1]) + a++; +// CHECK-NEXT: a++; + #pragma omp target map(this[a]) +// CHECK-NEXT: #pragma omp target map(tofrom: this[a]) + a++; +// CHECK-NEXT: a++; + } +// CHECK-NEXT: } +}; +// CHECK-NEXT: }; + // CHECK-LABEL: int main(int argc, char **argv) { int main (int argc, char **argv) { int i, j, a[20], always, close; diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp index 6395dd354fa..a5026cf660b 100644 --- a/clang/test/OpenMP/target_codegen.cpp +++ b/clang/test/OpenMP/target_codegen.cpp @@ -40,6 +40,7 @@ // CHECK-DAG: [[TT:%.+]] = type { i64, i8 } // CHECK-DAG: [[S1:%.+]] = type { double } +// CHECK-DAG: [[S2:%.+]] = type { i32, i32, i32 } // CHECK-DAG: [[ENTTY:%.+]] = type { i8*, i8*, i[[SZ:32|64]], i32, i32 } // CHECK-DAG: [[DEVTY:%.+]] = type { i8*, i8*, [[ENTTY]]*, [[ENTTY]]* } // CHECK-DAG: [[DSCTY:%.+]] = type { i32, [[DEVTY]]*, [[ENTTY]]*, [[ENTTY]]* } @@ -48,8 +49,8 @@ // CHECK-DAG: $[[REGFN:\.omp_offloading\..+]] = comdat -// We have 8 target regions, but only 7 that actually will generate offloading -// code and have mapped arguments, and only 5 have all-constant map sizes. +// We have 9 target regions, but only 8 that actually will generate offloading +// code and have mapped arguments, and only 6 have all-constant map sizes. // CHECK-DAG: [[SIZET:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 0, i[[SZ]] 4] // CHECK-DAG: [[MAPT:@.+]] = private unnamed_addr constant [2 x i64] [i64 544, i64 800] @@ -63,6 +64,9 @@ // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40] // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547] // CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 288, i64 288, i64 547] +// CHECK-DAG: [[SIZET9:@.+]] = private unnamed_addr constant [1 x i[[SZ]]] [i[[SZ]] 12] +// CHECK-DAG: [[MAPT10:@.+]] = private unnamed_addr constant [1 x i64] [i64 35] +// CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 @@ -80,6 +84,7 @@ // TCHECK: @{{.+}} = weak constant [[ENTTY]] // TCHECK: @{{.+}} = weak constant [[ENTTY]] // TCHECK: @{{.+}} = weak constant [[ENTTY]] +// TCHECK: @{{.+}} = weak constant [[ENTTY]] // TCHECK-NOT: @{{.+}} = weak constant [[ENTTY]] // Check if offloading descriptor is created. @@ -691,6 +696,31 @@ int bar(int n){ // CHECK: [[IFEND]] +// CHECK: define {{.*}}@{{.*}}zee{{.*}} + +// CHECK: [[LOCAL_THIS:%.+]] = alloca [[S2]]* +// CHECK: [[BP:%.+]] = alloca [1 x i8*] +// CHECK: [[P:%.+]] = alloca [1 x i8*] +// CHECK: [[LOCAL_THIS1:%.+]] = load [[S2]]*, [[S2]]** [[LOCAL_THIS]] +// CHECK: [[ARR_IDX:%.+]] = getelementptr inbounds [[S2]], [[S2]]* [[LOCAL_THIS1]], i[[SZ]] 0 +// CHECK: [[ARR_IDX2:%.+]] = getelementptr inbounds [[S2]], [[S2]]* [[LOCAL_THIS1]], i[[SZ]] 0 + +// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0 +// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0 +// CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to [[S2]]** +// CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to [[S2]]** +// CHECK-DAG: store [[S2]]* [[ARR_IDX]], [[S2]]** [[CBPADDR0]] +// CHECK-DAG: store [[S2]]* [[ARR_IDX2]], [[S2]]** [[CPADDR0]] + +// CHECK: [[BPR:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0 +// CHECK: [[PR:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0 +// CHECK: [[RET:%.+]] = call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 1, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i[[SZ]]* getelementptr inbounds ([1 x i[[SZ]]], [1 x i[[SZ]]]* [[SIZET9]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPT10]], i32 0, i32 0)) +// CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 +// CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] +// CHECK: [[FAIL]] +// CHECK: call void [[HVT0:@.+]]([[S2]]* [[LOCAL_THIS1]]) +// CHECK-NEXT: br label %[[END]] +// CHECK: [[END]] // Check that the offloading functions are emitted and that the arguments are // correct and loaded correctly for the target regions of the callees of bar(). @@ -765,4 +795,20 @@ void bar () { pragma_target {} } + +class S2 { + int a, b, c; + +public: + void zee() { + #pragma omp target map(this[0]) + a++; + } +}; + +int main () { + S2 bar; + bar.zee(); +} + #endif diff --git a/clang/test/OpenMP/target_messages.cpp b/clang/test/OpenMP/target_messages.cpp index 4fa8272ab45..9bd8b3749e0 100644 --- a/clang/test/OpenMP/target_messages.cpp +++ b/clang/test/OpenMP/target_messages.cpp @@ -43,6 +43,18 @@ void bar() { void foo() { } +class S { + public: + void zee() { + #pragma omp target map(this[:2]) // expected-note {{expected length on mapping of 'this' array section expression to be '1'}} // expected-error {{invalid 'this' expression on 'map' clause}} + int a; + #pragma omp target map(this[1:1]) // expected-note {{expected lower bound on mapping of 'this' array section expression to be '0' or not specified}} // expected-error {{invalid 'this' expression on 'map' clause}} + int b; + #pragma omp target map(this[1]) // expected-note {{expected 'this' subscript expression on map clause to be 'this[0]'}} // expected-error {{invalid 'this' expression on 'map' clause}} + int c; + } +}; + #pragma omp target // expected-error {{unexpected OpenMP directive '#pragma omp target'}} int main(int argc, char **argv) { |

