r350252 - [OpenMP] Added support for explicit mapping of classes using 'this' pointer. Differential revision: https://reviews.llvm.org/D55982
Patrick Lyster via cfe-commits
cfe-commits at lists.llvm.org
Wed Jan 2 11:28:48 PST 2019
Author: plyster
Date: Wed Jan 2 11:28:48 2019
New Revision: 350252
URL: http://llvm.org/viewvc/llvm-project?rev=350252&view=rev
Log:
[OpenMP] Added support for explicit mapping of classes using 'this' pointer. Differential revision: https://reviews.llvm.org/D55982
Modified:
cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/test/OpenMP/target_ast_print.cpp
cfe/trunk/test/OpenMP/target_codegen.cpp
cfe/trunk/test/OpenMP/target_messages.cpp
Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=350252&r1=350251&r2=350252&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
+++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Wed Jan 2 11:28:48 2019
@@ -9076,6 +9076,14 @@ def note_omp_requires_previous_clause :
"%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 {
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=350252&r1=350251&r2=350252&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Wed Jan 2 11:28:48 2019
@@ -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 =
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=350252&r1=350251&r2=350252&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Wed Jan 2 11:28:48 2019
@@ -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 *checkMapClauseExpress
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 *checkMapClauseExpress
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 &SemaRe
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).
Modified: cfe/trunk/test/OpenMP/target_ast_print.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_ast_print.cpp?rev=350252&r1=350251&r2=350252&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_ast_print.cpp (original)
+++ cfe/trunk/test/OpenMP/target_ast_print.cpp Wed Jan 2 11:28:48 2019
@@ -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;
Modified: cfe/trunk/test/OpenMP/target_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_codegen.cpp?rev=350252&r1=350251&r2=350252&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_codegen.cpp Wed Jan 2 11:28:48 2019
@@ -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
Modified: cfe/trunk/test/OpenMP/target_messages.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_messages.cpp?rev=350252&r1=350251&r2=350252&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_messages.cpp (original)
+++ cfe/trunk/test/OpenMP/target_messages.cpp Wed Jan 2 11:28:48 2019
@@ -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) {
More information about the cfe-commits
mailing list