jdenny created this revision.
jdenny added reviewers: ABataev, jdoerfert, hfinkel, kkwli0.
Herald added a subscriber: guansong.
Herald added a project: clang.
Without this patch, each of the following `map` clauses doesn't map
its variable into the target region because the variable is unused in
the target region, as discussed in D65835#1624669
<https://reviews.llvm.org/D65835#1624669>:
#pragma omp target map(a)
{}
#pragma omp target map(a)
#pragma omp teams private(a)
{
a++;
}
This patch fixes that by marking all map clause variables for
capturing. That means the capturing analysis now sometimes runs on a
capture region within a combined construct, so this patch adjusts the
analysis to be precise about how many capture regions remain in a
combined construct. Otherwise, existing tests break.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D66247
Files:
clang/include/clang/Sema/ScopeInfo.h
clang/include/clang/Sema/Sema.h
clang/lib/Sema/Sema.cpp
clang/lib/Sema/SemaExpr.cpp
clang/lib/Sema/SemaOpenMP.cpp
clang/test/OpenMP/target_map_codegen.cpp
Index: clang/test/OpenMP/target_map_codegen.cpp
===================================================================
--- clang/test/OpenMP/target_map_codegen.cpp
+++ clang/test/OpenMP/target_map_codegen.cpp
@@ -5329,5 +5329,125 @@
// CK31: define {{.+}}[[CALL00]]
// CK31: define {{.+}}[[CALL01]]
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK32 --check-prefix CK32-64
+// RUN: %clang_cc1 -DCK32 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK32 --check-prefix CK32-64
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK32 --check-prefix CK32-32
+// RUN: %clang_cc1 -DCK32 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK32 --check-prefix CK32-32
+
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK32 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s
+// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
+#ifdef CK32
+
+// CK32-LABEL: @.__omp_offloading_{{.*}}map_unused_var{{.*}}_l5383.region_id = weak constant i8 0
+// CK32: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK32: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+
+// CK32-LABEL: @.__omp_offloading_{{.*}}map_unused_var{{.*}}_l5399.region_id = weak constant i8 0
+// CK32: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK32: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
+
+// CK32-LABEL: @.__omp_offloading_{{.*}}map_unused_var{{.*}}_l5416.region_id = weak constant i8 0
+// CK32: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 396]
+// CK32: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+
+// CK32-LABEL: @.__omp_offloading_{{.*}}map_unused_var{{.*}}_l5432.region_id = weak constant i8 0
+// CK32: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+// CK32: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
+
+// CK32-LABEL: map_unused_var{{.*}}(
+void map_unused_var (){
+ float a;
+
+ // Region 00: default map type
+ // CK32-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+ // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK32-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+ // CK32-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+ // CK32-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+ // CK32-DAG: store float* [[VAR0]], float** [[CP0]]
+
+ // CK32: call void [[CALL00:@.+]](float* {{[^,]+}})
+ #pragma omp target map(a)
+ {}
+
+ // Region 01: non-default map type
+ // CK32-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+ // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK32-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float**
+ // CK32-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float**
+ // CK32-DAG: store float* [[VAR1:%.+]], float** [[CBP1]]
+ // CK32-DAG: store float* [[VAR1]], float** [[CP1]]
+
+ // CK32: call void [[CALL01:@.+]](float* {{[^,]+}})
+ #pragma omp target map(to: a)
+ {}
+
+ // Region 02: non-scalar data type
+ // CK32-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}})
+ // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK32-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [99 x float]**
+ // CK32-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to [99 x float]**
+ // CK32-DAG: store [99 x float]* [[VAR2:%.+]], [99 x float]** [[CBP2]]
+ // CK32-DAG: store [99 x float]* [[VAR2]], [99 x float]** [[CP2]]
+
+ // CK32: call void [[CALL02:@.+]]([99 x float]* {{[^,]+}})
+ float arr[99];
+ #pragma omp target map(arr)
+ {}
+
+ // Region 03: used, but only in nested private region
+ // CK32-DAG: call i32 @__tgt_target_teams(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i32 0, i32 0)
+ // CK32-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+ // CK32-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+ // CK32-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+ // CK32-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to float**
+ // CK32-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to float**
+ // CK32-DAG: store float* [[VAR3:%.+]], float** [[CBP3]]
+ // CK32-DAG: store float* [[VAR3]], float** [[CP3]]
+
+ // CK32: call void [[CALL03:@.+]](float* {{[^,]+}})
+ #pragma omp target map(a)
+ #pragma omp teams private(a)
+ {
+ a++;
+ }
+}
+// CK32: define {{.+}}[[CALL00]]
+// CK32-NOT: call {{.*\.omp_outlined\.}}
+
+// CK32: define {{.+}}[[CALL01]]
+// CK32-NOT: call {{.*\.omp_outlined\.}}
+
+// CK32: define {{.+}}[[CALL02]]
+// CK32-NOT: call {{.*\.omp_outlined\.}}
+
+// CK32: define {{.+}}[[CALL03]]
+// CK32: call {{.*}} [[OUTLINE03:@\.omp_outlined\.[^ ]*]]
+// CK32: define {{.+}}[[OUTLINE03]]
+// CK32: alloca float
+
#endif
#endif
Index: clang/lib/Sema/SemaOpenMP.cpp
===================================================================
--- clang/lib/Sema/SemaOpenMP.cpp
+++ clang/lib/Sema/SemaOpenMP.cpp
@@ -1853,13 +1853,6 @@
return nullptr;
}
-void Sema::adjustOpenMPTargetScopeIndex(unsigned &FunctionScopesIndex,
- unsigned Level) const {
- SmallVector<OpenMPDirectiveKind, 4> Regions;
- getOpenMPCaptureRegions(Regions, DSAStack->getDirective(Level));
- FunctionScopesIndex -= Regions.size();
-}
-
void Sema::startOpenMPLoop() {
assert(LangOpts.OpenMP && "OpenMP must be enabled.");
if (isOpenMPLoopDirective(DSAStack->getCurrentDirective()))
@@ -3408,6 +3401,7 @@
OMPScheduleClause *SC = nullptr;
SmallVector<const OMPLinearClause *, 4> LCs;
SmallVector<const OMPClauseWithPreInit *, 4> PICs;
+ SmallVector<OMPMapClause *, 4> MCs;
// This is required for proper codegen.
for (OMPClause *Clause : Clauses) {
if (isOpenMPTaskingDirective(DSAStack->getCurrentDirective()) &&
@@ -3447,6 +3441,8 @@
OC = cast<OMPOrderedClause>(Clause);
else if (Clause->getClauseKind() == OMPC_linear)
LCs.push_back(cast<OMPLinearClause>(Clause));
+ else if (Clause->getClauseKind() == OMPC_map)
+ MCs.push_back(cast<OMPMapClause>(Clause));
}
// OpenMP, 2.7.1 Loop Construct, Restrictions
// The nonmonotonic modifier cannot be specified if an ordered clause is
@@ -3503,6 +3499,14 @@
}
}
}
+ if (ThisCaptureRegion == OMPD_target) {
+ for (OMPMapClause *MC : MCs) {
+ for (ValueDecl *D : MC->all_decls()) {
+ if (auto *VD = dyn_cast_or_null<VarDecl>(D))
+ MarkVariableReferenced(VD->getLocation(), VD);
+ }
+ }
+ }
if (++CompletedRegions == CaptureRegions.size())
DSAStack->setBodyComplete();
SR = ActOnCapturedRegionEnd(SR.get());
Index: clang/lib/Sema/SemaExpr.cpp
===================================================================
--- clang/lib/Sema/SemaExpr.cpp
+++ clang/lib/Sema/SemaExpr.cpp
@@ -15784,7 +15784,7 @@
// target region, therefore we need to propagate the capture from the
// enclosing region. Therefore, the capture is not initially nested.
if (IsTargetCap)
- adjustOpenMPTargetScopeIndex(FunctionScopesIndex, RSI->OpenMPLevel);
+ FunctionScopesIndex -= RSI->OpenMPCaptureLevel + 1;
if (IsTargetCap || IsOpenMPPrivateDecl) {
Nested = !IsTargetCap;
Index: clang/lib/Sema/Sema.cpp
===================================================================
--- clang/lib/Sema/Sema.cpp
+++ clang/lib/Sema/Sema.cpp
@@ -2108,10 +2108,16 @@
void Sema::PushCapturedRegionScope(Scope *S, CapturedDecl *CD, RecordDecl *RD,
CapturedRegionKind K) {
- CapturingScopeInfo *CSI = new CapturedRegionScopeInfo(
+ CapturedRegionScopeInfo *CSI = new CapturedRegionScopeInfo(
getDiagnostics(), S, CD, RD, CD->getContextParam(), K,
(getLangOpts().OpenMP && K == CR_OpenMP) ? getOpenMPNestingLevel() : 0);
CSI->ReturnType = Context.VoidTy;
+ if (getLangOpts().OpenMP && K == CR_OpenMP) {
+ if (auto *P = dyn_cast<CapturedRegionScopeInfo>(FunctionScopes.back())) {
+ if (P->CapRegionKind == CR_OpenMP && CSI->OpenMPLevel == P->OpenMPLevel)
+ CSI->OpenMPCaptureLevel = P->OpenMPCaptureLevel + 1;
+ }
+ }
FunctionScopes.push_back(CSI);
}
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -8980,10 +8980,6 @@
/// Returns OpenMP nesting level for current directive.
unsigned getOpenMPNestingLevel() const;
- /// Adjusts the function scopes index for the target-based regions.
- void adjustOpenMPTargetScopeIndex(unsigned &FunctionScopesIndex,
- unsigned Level) const;
-
/// Push new OpenMP function region for non-capturing function.
void pushOpenMPFunctionRegion();
Index: clang/include/clang/Sema/ScopeInfo.h
===================================================================
--- clang/include/clang/Sema/ScopeInfo.h
+++ clang/include/clang/Sema/ScopeInfo.h
@@ -756,13 +756,15 @@
unsigned short CapRegionKind;
unsigned short OpenMPLevel;
+ unsigned short OpenMPCaptureLevel;
CapturedRegionScopeInfo(DiagnosticsEngine &Diag, Scope *S, CapturedDecl *CD,
RecordDecl *RD, ImplicitParamDecl *Context,
CapturedRegionKind K, unsigned OpenMPLevel)
: CapturingScopeInfo(Diag, ImpCap_CapturedRegion),
TheCapturedDecl(CD), TheRecordDecl(RD), TheScope(S),
- ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel) {
+ ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel),
+ OpenMPCaptureLevel(0) {
Kind = SK_CapturedRegion;
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
