sfantao created this revision.
sfantao added reviewers: ABataev, hfinkel, carlo.bertolli, arpith-jacob, kkwli0.
sfantao added subscribers: fraggamuffin, caomhin, cfe-commits.
Target regions require globals to be captured. The capturing scope has been
determine by matching the scope of the captured region and the data sharing
attribute stack (DSA stack). However, if the code is in a templated function,
the scope is set to null and can no longer be used to determine scopes.
This patch fixes the issue by matching the SEMA parent contexts (that are
always defined) of the region scope and DSA scope.
http://reviews.llvm.org/D18110
Files:
lib/Sema/SemaOpenMP.cpp
test/OpenMP/target_codegen_global_capture.cpp
test/OpenMP/teams_codegen.cpp
Index: test/OpenMP/teams_codegen.cpp
===================================================================
--- test/OpenMP/teams_codegen.cpp
+++ test/OpenMP/teams_codegen.cpp
@@ -250,12 +250,10 @@
// CK4: ret void
// CK4-NEXT: }
-// CK4: define {{.*}}void @{{[^,]+}}(i8*** dereferenceable({{.}}) [[ARGC1:%.+]])
-// CK4: [[ARGCADDR1:%.+]] = alloca i8***
-// CK4: store i8*** [[ARGC1]], i8**** [[ARGCADDR1]]
-// CK4: [[CONV1:%.+]] = load i8***, i8**** [[ARGCADDR1]]
-// CK4: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* {{.+}} to void (i32*, i32*, ...)*), i8*** [[CONV1]])
-
+// CK4: define {{.*}}void @{{[^,]+}}(i8** [[ARGC1:%.+]])
+// CK4: [[ARGCADDR1:%.+]] = alloca i8**
+// CK4: store i8** [[ARGC1]], i8*** [[ARGCADDR1]]
+// CK4: call {{.*}}void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i8***)* {{.+}} to void (i32*, i32*, ...)*), i8*** [[ARGCADDR1]])
#endif // CK4
@@ -320,21 +318,22 @@
// CK5-64: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[CONV]])
// CK5-32: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*), i32* [[ARGCADDR]])
-// CK5: define {{.*}}void @{{[^,]+}}(i{{.+}} dereferenceable({{.+}}) [[AP:%.+]], i{{.+}} dereferenceable({{.+}}) [[BP:%.+]], i{{.+}} dereferenceable({{.+}}) [[ARGC:%.+]])
+// CK5: define {{.*}}void @{{[^,]+}}(i{{.+}} [[AP:%.+]], i{{.+}} [[BP:%.+]], i8** [[ARGC:%.+]])
// CK5: [[AADDR:%.+]] = alloca i{{.+}}
// CK5: [[BADDR:%.+]] = alloca i{{.+}}
-// CK5: [[ARGCADDR:%.+]] = alloca i{{.+}}
+// CK5: [[ARGCADDR:%.+]] = alloca i8**
// CK5: [[GBL_TH_NUM:%.+]] = call i32 @__kmpc_global_thread_num(%ident_t* [[DEF_LOC_0]])
// CK5: store i{{.+}} [[AP]], i{{.+}}* [[AADDR]]
// CK5: store i{{.+}} [[BP]], i{{.+}}* [[BADDR]]
-// CK5: store i{{.+}} [[ARGC]], i{{.+}}* [[ARGCADDR]]
-// CK5: [[A_ADDR_VAL:%.+]] = load i32*, i32** [[AADDR]]
-// CK5: [[B_ADDR_VAL:%.+]] = load i32*, i32** [[BADDR]]
-// CK5: [[ARGC_ADDR_VAL:%.+]] = load i{{.+}}, i{{.+}}* [[ARGCADDR]]
-// CK5: [[A_VAL:%.+]] = load i32, i32* [[A_ADDR_VAL]]
-// CK5: [[B_VAL:%.+]] = load i32, i32* [[B_ADDR_VAL]]
-// CK5: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[A_VAL]], i32 [[B_VAL]])
-// CK5: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i{{.+}})* @.omp_outlined.{{.+}} to void (i32*, i32*, ...)*), i{{.+}} [[ARGC_ADDR_VAL]])
+// CK5: store i8** [[ARGC]], i8*** [[ARGCADDR]]
+// CK5-64: [[ACONV:%.+]] = bitcast i64* [[AADDR]] to i32*
+// CK5-64: [[BCONV:%.+]] = bitcast i64* [[BADDR]] to i32*
+// CK5-64: [[ACONVVAL:%.+]] = load i32, i32* [[ACONV]]
+// CK5-64: [[BCONVVAL:%.+]] = load i32, i32* [[BCONV]]
+// CK5-32: [[ACONVVAL:%.+]] = load i32, i32* [[AADDR]]
+// CK5-32: [[BCONVVAL:%.+]] = load i32, i32* [[BADDR]]
+// CK5: {{.+}} = call i32 @__kmpc_push_num_teams(%ident_t* [[DEF_LOC_0]], i32 [[GBL_TH_NUM]], i32 [[ACONVVAL]], i32 [[BCONVVAL]])
+// CK5: call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%ident_t* [[DEF_LOC_0]], i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i{{.+}})* @.omp_outlined.{{.+}} to void (i32*, i32*, ...)*), i8*** [[ARGCADDR]])
// CK5: ret void
// CK5-NEXT: }
Index: test/OpenMP/target_codegen_global_capture.cpp
===================================================================
--- test/OpenMP/target_codegen_global_capture.cpp
+++ test/OpenMP/target_codegen_global_capture.cpp
@@ -21,6 +21,11 @@
// CHECK-DAG: [[BB:@.+]] = internal global float 1.000000e+01
// CHECK-DAG: [[BC:@.+]] = internal global float 1.100000e+01
// CHECK-DAG: [[BD:@.+]] = internal global float 1.200000e+01
+// CHECK-DAG: [[TBA:@.+]] = {{.*}}global float 1.700000e+01
+// CHECK-DAG: [[TBB:@.+]] = {{.*}}global float 1.800000e+01
+// CHECK-DAG: [[TBC:@.+]] = {{.*}}global float 1.900000e+01
+// CHECK-DAG: [[TBD:@.+]] = {{.*}}global float 2.000000e+01
+
double Ga = 1.0;
double Gb = 2.0;
double Gc = 3.0;
@@ -42,14 +47,14 @@
static float Sd = 8.0;
// CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LB]],
- // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb,
- // CHECK-DAG: [[VALFB:%.+]] = load float, float* @_ZZ3foossssE2Sb,
- // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc,
+ // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
+ // CHECK-DAG: [[VALFB:%.+]] = load float, float* [[FB]],
+ // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
// CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LC]],
- // CHECK-DAG: [[VALFC:%.+]] = load float, float* @_ZZ3foossssE2Sc,
+ // CHECK-DAG: [[VALFC:%.+]] = load float, float* [[FC]],
// CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LD]],
- // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd,
- // CHECK-DAG: [[VALFD:%.+]] = load float, float* @_ZZ3foossssE2Sd,
+ // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
+ // CHECK-DAG: [[VALFD:%.+]] = load float, float* [[FD]],
// 3 local vars being captured.
@@ -178,14 +183,156 @@
#pragma omp parallel
{
// CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LLB]],
- // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* @Gb,
- // CHECK-DAG: [[VALFB:%.+]] = load float, float* @_ZZ3barssssE2Sb,
- // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* @Gc,
+ // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
+ // CHECK-DAG: [[VALFB:%.+]] = load float, float* [[BB]],
+ // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
+ // CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LLC]],
+ // CHECK-DAG: [[VALFC:%.+]] = load float, float* [[BC]],
+ // CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LLD]],
+ // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
+ // CHECK-DAG: [[VALFD:%.+]] = load float, float* [[BD]],
+
+ // 3 local vars being captured.
+
+ // CHECK-DAG: store i16 [[VALLB]], i16* [[CONVLB:%.+]],
+ // CHECK-DAG: [[CONVLB]] = bitcast i[[sz:64|32]]* [[CADDRLB:%.+]] to i16*
+ // CHECK-DAG: [[CVALLB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLB]],
+ // CHECK-DAG: [[CPTRLB:%.+]] = inttoptr i[[sz]] [[CVALLB]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLB]], i8** [[GEPLB:%.+]],
+ // CHECK-DAG: [[GEPLB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store i16 [[VALLC]], i16* [[CONVLC:%.+]],
+ // CHECK-DAG: [[CONVLC]] = bitcast i[[sz]]* [[CADDRLC:%.+]] to i16*
+ // CHECK-DAG: [[CVALLC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLC]],
+ // CHECK-DAG: [[CPTRLC:%.+]] = inttoptr i[[sz]] [[CVALLC]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLC]], i8** [[GEPLC:%.+]],
+ // CHECK-DAG: [[GEPLC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store i16 [[VALLD]], i16* [[CONVLD:%.+]],
+ // CHECK-DAG: [[CONVLD]] = bitcast i[[sz]]* [[CADDRLD:%.+]] to i16*
+ // CHECK-DAG: [[CVALLD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRLD]],
+ // CHECK-DAG: [[CPTRLD:%.+]] = inttoptr i[[sz]] [[CVALLD]] to i8*
+ // CHECK-DAG: store i8* [[CPTRLD]], i8** [[GEPLD:%.+]],
+ // CHECK-DAG: [[GEPLD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // 3 static vars being captured.
+
+ // CHECK-DAG: store float [[VALFB]], float* [[CONVFB:%.+]],
+ // CHECK-DAG: [[CONVFB]] = bitcast i[[sz]]* [[CADDRFB:%.+]] to float*
+ // CHECK-DAG: [[CVALFB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFB]],
+ // CHECK-DAG: [[CPTRFB:%.+]] = inttoptr i[[sz]] [[CVALFB]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFB]], i8** [[GEPFB:%.+]],
+ // CHECK-DAG: [[GEPFB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store float [[VALFC]], float* [[CONVFC:%.+]],
+ // CHECK-DAG: [[CONVFC]] = bitcast i[[sz]]* [[CADDRFC:%.+]] to float*
+ // CHECK-DAG: [[CVALFC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFC]],
+ // CHECK-DAG: [[CPTRFC:%.+]] = inttoptr i[[sz]] [[CVALFC]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFC]], i8** [[GEPFC:%.+]],
+ // CHECK-DAG: [[GEPFC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-DAG: store float [[VALFD]], float* [[CONVFD:%.+]],
+ // CHECK-DAG: [[CONVFD]] = bitcast i[[sz]]* [[CADDRFD:%.+]] to float*
+ // CHECK-DAG: [[CVALFD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRFD]],
+ // CHECK-DAG: [[CPTRFD:%.+]] = inttoptr i[[sz]] [[CVALFD]] to i8*
+ // CHECK-DAG: store i8* [[CPTRFD]], i8** [[GEPFD:%.+]],
+ // CHECK-DAG: [[GEPFD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // 3 static global vars being captured.
+
+ // CHECK-64-DAG: store double [[VALGB]], double* [[CONVGB:%.+]],
+ // CHECK-64-DAG: [[CONVGB]] = bitcast i[[sz]]* [[CADDRGB:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGB:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGB]],
+ // CHECK-64-DAG: [[CPTRGB:%.+]] = inttoptr i[[sz]] [[CVALGB]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGB]], i8** [[GEPGB:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gb to i8*), i8** [[GEPGB:%.+]],
+ // CHECK-DAG: [[GEPGB]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-64-DAG: store double [[VALGC]], double* [[CONVGC:%.+]],
+ // CHECK-64-DAG: [[CONVGC]] = bitcast i[[sz]]* [[CADDRGC:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGC:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGC]],
+ // CHECK-64-DAG: [[CPTRGC:%.+]] = inttoptr i[[sz]] [[CVALGC]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGC]], i8** [[GEPGC:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gc to i8*), i8** [[GEPGC:%.+]],
+ // CHECK-DAG: [[GEPGC]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK-64-DAG: store double [[VALGD]], double* [[CONVGD:%.+]],
+ // CHECK-64-DAG: [[CONVGD]] = bitcast i[[sz]]* [[CADDRGD:%.+]] to double*
+ // CHECK-64-DAG: [[CVALGD:%.+]] = load i[[sz]], i[[sz]]* [[CADDRGD]],
+ // CHECK-64-DAG: [[CPTRGD:%.+]] = inttoptr i[[sz]] [[CVALGD]] to i8*
+ // CHECK-64-DAG: store i8* [[CPTRGD]], i8** [[GEPGD:%.+]],
+ // CHECK-32-DAG: store i8* bitcast (double* @Gd to i8*), i8** [[GEPGD:%.+]],
+ // CHECK-DAG: [[GEPGD]] = getelementptr inbounds [9 x i8*], [9 x i8*]* %{{.+}}, i32 0, i32 {{[0-8]}}
+
+ // CHECK: call i32 @__tgt_target
+ // CHECK: call void [[OFFLOADF:@.+]](
+ // Capture b, Gb, Sb, Gc, c, Sc, d, Gd, Sd
+ #pragma omp target if(Ga>0.0 && a>0 && Sa>0.0)
+ {
+ b += 1;
+ Gb += 1.0;
+ Sb += 1.0;
+
+ // CHECK: define internal void [[OFFLOADF]]({{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}}, {{.+}} {{.*}}%{{.+}})
+ // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}})
+
+ // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}}, {{.+}}* dereferenceable({{.+}}) %{{.+}})
+ // Capture d, Gd, Sd
+ #pragma omp parallel if(Gc>0.0 && c>0 && Sc>0.0)
+ {
+ d += 1;
+ Gd += 1.0;
+ Sd += 1.0;
+ }
+ }
+ }
+ return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
+}
+
+///
+/// Tests with template functions.
+///
+
+// CHECK: define {{.*}} @{{.*}}tbar2{{.*}}(
+
+// CHECK: define {{.*}} @{{.*}}tbar{{.*}}(
+// CHECK-SAME: i16 {{[^,]*}}[[A:%[^,]+]],
+// CHECK-SAME: i16 {{[^,]*}}[[B:%[^,]+]],
+// CHECK-SAME: i16 {{[^,]*}}[[C:%[^,]+]],
+// CHECK-SAME: i16 {{[^,]*}}[[D:%[^,]+]])
+// CHECK: [[LA:%.+]] = alloca i16
+// CHECK: [[LB:%.+]] = alloca i16
+// CHECK: [[LC:%.+]] = alloca i16
+// CHECK: [[LD:%.+]] = alloca i16
+template<typename T>
+int tbar(T a, T b, T c, T d){
+ static float Sa = 17.0;
+ static float Sb = 18.0;
+ static float Sc = 19.0;
+ static float Sd = 20.0;
+
+ // CHECK: call void {{.*}}@__kmpc_fork_call(%ident_t* {{.+}}, i32 {{.+}}, void (i32*, i32*, ...)* bitcast ({{.*}}[[PARF:@.+]] to {{.*}}), i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}}, i16* %{{.+}})
+ // CHECK: define internal void [[PARF]](i32* noalias %{{.*}}, i32* noalias %{{.*}}, i16* dereferenceable(2) [[A:%.+]], i16* dereferenceable(2) [[B:%.+]], i16* dereferenceable(2) [[C:%.+]], i16* dereferenceable(2) [[D:%.+]])
+ // Capture a, b, c, d
+ // CHECK: [[ALLOCLA:%.+]] = alloca i16
+ // CHECK: [[ALLOCLB:%.+]] = alloca i16
+ // CHECK: [[ALLOCLC:%.+]] = alloca i16
+ // CHECK: [[ALLOCLD:%.+]] = alloca i16
+ // CHECK: [[LLA:%.+]] = load i16*, i16** [[ALLOCLA]],
+ // CHECK: [[LLB:%.+]] = load i16*, i16** [[ALLOCLB]],
+ // CHECK: [[LLC:%.+]] = load i16*, i16** [[ALLOCLC]],
+ // CHECK: [[LLD:%.+]] = load i16*, i16** [[ALLOCLD]],
+ #pragma omp parallel
+ {
+ // CHECK-DAG: [[VALLB:%.+]] = load i16, i16* [[LLB]],
+ // CHECK-64-DAG: [[VALGB:%.+]] = load double, double* [[GB]],
+ // CHECK-DAG: [[VALFB:%.+]] = load float, float* [[TBB]],
+ // CHECK-64-DAG: [[VALGC:%.+]] = load double, double* [[GC]],
// CHECK-DAG: [[VALLC:%.+]] = load i16, i16* [[LLC]],
- // CHECK-DAG: [[VALFC:%.+]] = load float, float* @_ZZ3barssssE2Sc,
+ // CHECK-DAG: [[VALFC:%.+]] = load float, float* [[TBC]],
// CHECK-DAG: [[VALLD:%.+]] = load i16, i16* [[LLD]],
- // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* @Gd,
- // CHECK-DAG: [[VALFD:%.+]] = load float, float* @_ZZ3barssssE2Sd,
+ // CHECK-64-DAG: [[VALGD:%.+]] = load double, double* [[GD]],
+ // CHECK-DAG: [[VALFD:%.+]] = load float, float* [[TBD]],
// 3 local vars being captured.
@@ -284,4 +431,8 @@
return a + b + c + d + (int)Sa + (int)Sb + (int)Sc + (int)Sd;
}
+int tbar2(short a, short b, short c, short d){
+ return tbar(a, b, c, d);
+}
+
#endif
Index: lib/Sema/SemaOpenMP.cpp
===================================================================
--- lib/Sema/SemaOpenMP.cpp
+++ lib/Sema/SemaOpenMP.cpp
@@ -113,17 +113,21 @@
bool CancelRegion;
unsigned AssociatedLoops;
SourceLocation InnerTeamsRegionLoc;
+ const DeclContext *ParentDeclContext;
SharingMapTy(OpenMPDirectiveKind DKind, DeclarationNameInfo Name,
- Scope *CurScope, SourceLocation Loc)
+ Scope *CurScope, SourceLocation Loc,
+ const DeclContext *ParentDeclContext)
: SharingMap(), AlignedMap(), LCVMap(), DefaultAttr(DSA_unspecified),
Directive(DKind), DirectiveName(std::move(Name)), CurScope(CurScope),
ConstructLoc(Loc), OrderedRegion(), NowaitRegion(false),
- CancelRegion(false), AssociatedLoops(1), InnerTeamsRegionLoc() {}
+ CancelRegion(false), AssociatedLoops(1), InnerTeamsRegionLoc(),
+ ParentDeclContext(ParentDeclContext) {}
SharingMapTy()
: SharingMap(), AlignedMap(), LCVMap(), DefaultAttr(DSA_unspecified),
Directive(OMPD_unknown), DirectiveName(), CurScope(nullptr),
ConstructLoc(), OrderedRegion(), NowaitRegion(false),
- CancelRegion(false), AssociatedLoops(1), InnerTeamsRegionLoc() {}
+ CancelRegion(false), AssociatedLoops(1), InnerTeamsRegionLoc(),
+ ParentDeclContext(nullptr) {}
};
typedef SmallVector<SharingMapTy, 4> StackTy;
@@ -156,8 +160,10 @@
void setForceVarCapturing(bool V) { ForceCapturing = V; }
void push(OpenMPDirectiveKind DKind, const DeclarationNameInfo &DirName,
- Scope *CurScope, SourceLocation Loc) {
- Stack.push_back(SharingMapTy(DKind, DirName, CurScope, Loc));
+ Scope *CurScope, SourceLocation Loc,
+ const DeclContext *ParentDeclContext) {
+ Stack.push_back(
+ SharingMapTy(DKind, DirName, CurScope, Loc, ParentDeclContext));
Stack.back().DefaultAttrLoc = Loc;
}
@@ -245,8 +251,9 @@
return Stack[Stack.size() - 2].Directive;
return OMPD_unknown;
}
- /// \brief Return the directive associated with the provided scope.
- OpenMPDirectiveKind getDirectiveForScope(const Scope *S) const;
+ /// \brief Return the directive associated with the provided region scope.
+ OpenMPDirectiveKind
+ getDirectiveForRegionScope(const CapturedRegionScopeInfo *RSI) const;
/// \brief Set default data sharing attribute to none.
void setDefaultDSANone(SourceLocation Loc) {
@@ -785,6 +792,9 @@
template <class NamedDirectivesPredicate>
bool DSAStackTy::hasDirective(NamedDirectivesPredicate DPred, bool FromParent) {
+ // We look only in the enclosing region.
+ if (Stack.size() < 2)
+ return false;
auto StartI = std::next(Stack.rbegin());
auto EndI = std::prev(Stack.rend());
if (FromParent && StartI != EndI) {
@@ -797,9 +807,16 @@
return false;
}
-OpenMPDirectiveKind DSAStackTy::getDirectiveForScope(const Scope *S) const {
+OpenMPDirectiveKind DSAStackTy::getDirectiveForRegionScope(
+ const CapturedRegionScopeInfo *RSI) const {
+ // Directives are expected to have a context associated.
+ if (!(RSI && RSI->TheCapturedDecl))
+ return OMPD_unknown;
+
+ // A DSA refers to this captured region if the parent contexts match.
+ auto *ParentContext = RSI->TheCapturedDecl->getParent();
for (auto I = Stack.rbegin(), EE = Stack.rend(); I != EE; ++I)
- if (I->CurScope == S)
+ if (I->ParentDeclContext == ParentContext)
return I->Directive;
return OMPD_unknown;
}
@@ -818,7 +835,7 @@
bool IsByRef = true;
// Find the directive that is associated with the provided scope.
- auto DKind = DSAStack->getDirectiveForScope(RSI->TheScope);
+ auto DKind = DSAStack->getDirectiveForRegionScope(RSI);
auto Ty = D->getType();
if (isOpenMPTargetExecutionDirective(DKind)) {
@@ -914,8 +931,7 @@
if (DSAStack->getCurrentDirective() == OMPD_target &&
!DSAStack->isClauseParsingMode())
return VD;
- if (DSAStack->getCurScope() &&
- DSAStack->hasDirective(
+ if (DSAStack->hasDirective(
[](OpenMPDirectiveKind K, const DeclarationNameInfo &DNI,
SourceLocation Loc) -> bool {
return isOpenMPTargetExecutionDirective(K);
@@ -964,7 +980,9 @@
void Sema::StartOpenMPDSABlock(OpenMPDirectiveKind DKind,
const DeclarationNameInfo &DirName,
Scope *CurScope, SourceLocation Loc) {
- DSAStack->push(DKind, DirName, CurScope, Loc);
+ // The current Sema declarative context is the parent of the captured region
+ // that refers to this DSA block.
+ DSAStack->push(DKind, DirName, CurScope, Loc, CurContext);
PushExpressionEvaluationContext(PotentiallyEvaluated);
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
