summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--clang/lib/Sema/SemaOpenMP.cpp21
-rw-r--r--clang/test/OpenMP/distribute_simd_loop_messages.cpp16
-rw-r--r--clang/test/OpenMP/teams_distribute_codegen.cpp4
-rw-r--r--clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp2
-rw-r--r--clang/test/OpenMP/teams_distribute_simd_codegen.cpp2
5 files changed, 38 insertions, 7 deletions
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
index 8b1fca89dee..b3b8fd655f1 100644
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -2768,6 +2768,7 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
DSAStackTy *Stack;
Sema &SemaRef;
bool ErrorFound = false;
+ bool TryCaptureCXXThisMembers = false;
CapturedStmt *CS = nullptr;
llvm::SmallVector<Expr *, 4> ImplicitFirstprivate;
llvm::SmallVector<Expr *, 4> ImplicitMap;
@@ -2779,12 +2780,26 @@ class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
if (!S->hasAssociatedStmt() || !S->getAssociatedStmt())
return;
visitSubCaptures(S->getInnermostCapturedStmt());
+ // Try to capture inner this->member references to generate correct mappings
+ // and diagnostics.
+ if (TryCaptureCXXThisMembers ||
+ (isOpenMPTargetExecutionDirective(Stack->getCurrentDirective()) &&
+ llvm::any_of(S->getInnermostCapturedStmt()->captures(),
+ [](const CapturedStmt::Capture &C) {
+ return C.capturesThis();
+ }))) {
+ bool SavedTryCaptureCXXThisMembers = TryCaptureCXXThisMembers;
+ TryCaptureCXXThisMembers = true;
+ Visit(S->getInnermostCapturedStmt()->getCapturedStmt());
+ TryCaptureCXXThisMembers = SavedTryCaptureCXXThisMembers;
+ }
}
public:
void VisitDeclRefExpr(DeclRefExpr *E) {
- if (E->isTypeDependent() || E->isValueDependent() ||
- E->containsUnexpandedParameterPack() || E->isInstantiationDependent())
+ if (TryCaptureCXXThisMembers || E->isTypeDependent() ||
+ E->isValueDependent() || E->containsUnexpandedParameterPack() ||
+ E->isInstantiationDependent())
return;
if (auto *VD = dyn_cast<VarDecl>(E->getDecl())) {
// Check the datasharing rules for the expressions in the clauses.
@@ -3018,7 +3033,7 @@ public:
})) {
Visit(E->getBase());
}
- } else {
+ } else if (!TryCaptureCXXThisMembers) {
Visit(E->getBase());
}
}
diff --git a/clang/test/OpenMP/distribute_simd_loop_messages.cpp b/clang/test/OpenMP/distribute_simd_loop_messages.cpp
index 423e5f1b116..183fdfd90cf 100644
--- a/clang/test/OpenMP/distribute_simd_loop_messages.cpp
+++ b/clang/test/OpenMP/distribute_simd_loop_messages.cpp
@@ -4,6 +4,22 @@
// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -x c++ -std=c++11 -fexceptions -fcxx-exceptions -verify=expected,omp4 %s -Wuninitialized
// RUN: %clang_cc1 -fsyntax-only -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -fexceptions -fcxx-exceptions -verify=expected,omp5 %s -Wuninitialized
+class S5 {
+ int a;
+ S5() : a(0) {}
+
+public:
+ S5(int v) : a(v) {}
+ S5 &operator=(S5 &s) {
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute simd
+ for (int k = 0; k < s.a; ++k) // expected-warning {{Non-trivial type 'S5' is mapped, only trivial types are guaranteed to be mapped correctly}}
+ ++s.a;
+ return *this;
+ }
+};
+
static int sii;
// expected-note@+1 {{defined as threadprivate or thread local}}
#pragma omp threadprivate(sii)
diff --git a/clang/test/OpenMP/teams_distribute_codegen.cpp b/clang/test/OpenMP/teams_distribute_codegen.cpp
index 0065c6cbd51..151fa2da49f 100644
--- a/clang/test/OpenMP/teams_distribute_codegen.cpp
+++ b/clang/test/OpenMP/teams_distribute_codegen.cpp
@@ -158,7 +158,7 @@ struct SS{
// CK3: define {{.*}}i32 @{{.+}}foo{{.+}}(
int foo(void) {
- // CK3: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0)
+ // CK3: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* %{{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0)
// CK3: call void @[[OFFL1:.+]]([[SSI]]* %{{.+}})
#pragma omp target
#pragma omp teams distribute
@@ -174,7 +174,7 @@ struct SS{
// CK3: define internal void @[[OUTL1]]({{.+}})
// CK3: call void @__kmpc_for_static_init_4(
// CK3: call void @__kmpc_for_static_fini(
- // CK3: ret void
+ // CK3: ret void
return a[0];
}
diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp
index bc42f740ad7..43d3500d196 100644
--- a/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp
+++ b/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp
@@ -161,7 +161,7 @@ struct SS{
// CK3: define {{.*}}i32 @{{.+}}foo{{.+}}(
int foo(void) {
- // CK3: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0)
+ // CK3: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* %{{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 0)
// CK3: call void @[[OFFL1:.+]]([[SSI]]* %{{.+}})
#pragma omp target
#pragma omp teams distribute parallel for
diff --git a/clang/test/OpenMP/teams_distribute_simd_codegen.cpp b/clang/test/OpenMP/teams_distribute_simd_codegen.cpp
index ff2a752cfa3..0f68328e67a 100644
--- a/clang/test/OpenMP/teams_distribute_simd_codegen.cpp
+++ b/clang/test/OpenMP/teams_distribute_simd_codegen.cpp
@@ -165,7 +165,7 @@ struct SS{
// CK3: define {{.*}}i32 @{{.+}}foo{{.+}}(
int foo(void) {
- // CK3: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 1)
+ // CK3: call i32 @__tgt_target_teams(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* %{{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i32 0, i32 1)
// CK3: call void @[[OFFL1:.+]]([[SSI]]* %{{.+}})
#pragma omp target
#pragma omp teams distribute simd
OpenPOWER on IntegriCloud