r314995 - [OPENMP] Fix mapping|privatization of implicitly captured variables.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 5 10:51:39 PDT 2017
Author: abataev
Date: Thu Oct 5 10:51:39 2017
New Revision: 314995
URL: http://llvm.org/viewvc/llvm-project?rev=314995&view=rev
Log:
[OPENMP] Fix mapping|privatization of implicitly captured variables.
If the `defaultmap(tofrom:scalar)` clause is specified, the scalars must
be mapped with 'tofrom' modifiers, otherwise they must be captured as
firstprivates.
Modified:
cfe/trunk/lib/Sema/SemaOpenMP.cpp
cfe/trunk/test/OpenMP/target_map_codegen.cpp
cfe/trunk/test/OpenMP/teams_distribute_firstprivate_codegen.cpp
Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=314995&r1=314994&r2=314995&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original)
+++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Thu Oct 5 10:51:39 2017
@@ -45,7 +45,13 @@ namespace {
enum DefaultDataSharingAttributes {
DSA_unspecified = 0, /// \brief Data sharing attribute not specified.
DSA_none = 1 << 0, /// \brief Default data sharing attribute 'none'.
- DSA_shared = 1 << 1 /// \brief Default data sharing attribute 'shared'.
+ DSA_shared = 1 << 1, /// \brief Default data sharing attribute 'shared'.
+};
+
+/// Attributes of the defaultmap clause.
+enum DefaultMapAttributes {
+ DMA_unspecified, /// Default mapping is not specified.
+ DMA_tofrom_scalar, /// Default mapping is 'tofrom:scalar'.
};
/// \brief Stack for tracking declarations used in OpenMP directives and
@@ -115,6 +121,8 @@ private:
LoopControlVariablesMapTy LCVMap;
DefaultDataSharingAttributes DefaultAttr = DSA_unspecified;
SourceLocation DefaultAttrLoc;
+ DefaultMapAttributes DefaultMapAttr = DMA_unspecified;
+ SourceLocation DefaultMapAttrLoc;
OpenMPDirectiveKind Directive = OMPD_unknown;
DeclarationNameInfo DirectiveName;
Scope *CurScope = nullptr;
@@ -341,6 +349,12 @@ public:
Stack.back().first.back().DefaultAttr = DSA_shared;
Stack.back().first.back().DefaultAttrLoc = Loc;
}
+ /// Set default data mapping attribute to 'tofrom:scalar'.
+ void setDefaultDMAToFromScalar(SourceLocation Loc) {
+ assert(!isStackEmpty());
+ Stack.back().first.back().DefaultMapAttr = DMA_tofrom_scalar;
+ Stack.back().first.back().DefaultMapAttrLoc = Loc;
+ }
DefaultDataSharingAttributes getDefaultDSA() const {
return isStackEmpty() ? DSA_unspecified
@@ -350,6 +364,17 @@ public:
return isStackEmpty() ? SourceLocation()
: Stack.back().first.back().DefaultAttrLoc;
}
+ DefaultMapAttributes getDefaultDMA() const {
+ return isStackEmpty() ? DMA_unspecified
+ : Stack.back().first.back().DefaultMapAttr;
+ }
+ DefaultMapAttributes getDefaultDMAAtLevel(unsigned Level) const {
+ return Stack.back().first[Level].DefaultMapAttr;
+ }
+ SourceLocation getDefaultDMALocation() const {
+ return isStackEmpty() ? SourceLocation()
+ : Stack.back().first.back().DefaultMapAttrLoc;
+ }
/// \brief Checks if the specified variable is a threadprivate.
bool isThreadPrivate(VarDecl *D) {
@@ -1242,7 +1267,8 @@ bool Sema::IsOpenMPCapturedByRef(ValueDe
IsByRef = !(Ty->isPointerType() && IsVariableAssociatedWithSection);
} else {
// By default, all the data that has a scalar type is mapped by copy.
- IsByRef = !Ty->isScalarType();
+ IsByRef = !Ty->isScalarType() ||
+ DSAStack->getDefaultDMAAtLevel(Level) == DMA_tofrom_scalar;
}
}
@@ -1804,7 +1830,7 @@ public:
if (auto *VD = dyn_cast<VarDecl>(E->getDecl())) {
VD = VD->getCanonicalDecl();
// Skip internally declared variables.
- if (VD->isLocalVarDecl() && !CS->capturesVariable(VD))
+ if (VD->hasLocalStorage() && !CS->capturesVariable(VD))
return;
auto DVar = Stack->getTopDSA(VD, false);
@@ -1848,20 +1874,16 @@ public:
MC.getAssociatedExpression()));
});
})) {
- bool CapturedByCopy = false;
+ bool IsFirstprivate = false;
// By default lambdas are captured as firstprivates.
if (const auto *RD =
VD->getType().getNonReferenceType()->getAsCXXRecordDecl())
- if (RD->isLambda())
- CapturedByCopy = true;
- CapturedByCopy =
- CapturedByCopy ||
- llvm::any_of(
- CS->captures(), [VD](const CapturedStmt::Capture &I) {
- return I.capturesVariableByCopy() &&
- I.getCapturedVar()->getCanonicalDecl() == VD;
- });
- if (CapturedByCopy)
+ IsFirstprivate = RD->isLambda();
+ IsFirstprivate =
+ IsFirstprivate ||
+ (VD->getType().getNonReferenceType()->isScalarType() &&
+ Stack->getDefaultDMA() != DMA_tofrom_scalar);
+ if (IsFirstprivate)
ImplicitFirstprivate.emplace_back(E);
else
ImplicitMap.emplace_back(E);
@@ -11905,6 +11927,7 @@ OMPClause *Sema::ActOnOpenMPDefaultmapCl
<< Value << getOpenMPClauseName(OMPC_defaultmap);
return nullptr;
}
+ DSAStack->setDefaultDMAToFromScalar(StartLoc);
return new (Context)
OMPDefaultmapClause(StartLoc, LParenLoc, MLoc, KindLoc, EndLoc, Kind, M);
Modified: cfe/trunk/test/OpenMP/target_map_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_map_codegen.cpp?rev=314995&r1=314994&r2=314995&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/target_map_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/target_map_codegen.cpp Thu Oct 5 10:51:39 2017
@@ -360,8 +360,8 @@ void implicit_maps_host_global (int a){
// CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288
// CK7-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
-// Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_IS_FIRST | OMP_MAP_IMPLICIT = 547
-// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547]
+// Map types: OMP_MAP_TO | OMP_MAP_PRIVATE_PTR | OMP_MAP_FIRST_REF = 161
+// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 161]
// CK7-LABEL: implicit_maps_double
void implicit_maps_double (int a){
@@ -562,7 +562,7 @@ void implicit_maps_double_complex (int a
// CK11-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]]
// CK11: call void [[KERNEL:@.+]]({ double, double }* [[PTR]])
- #pragma omp target
+ #pragma omp target defaultmap(tofrom:scalar)
{
dc *= dc;
}
@@ -589,8 +589,8 @@ void implicit_maps_double_complex (int a
// CK12-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8]
// Map types: OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288
// CK12-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288]
-// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547
-// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547]
+// Map types: OMP_MAP_TO | OMP_MAP_PRIVATE_PTR | OMP_MAP_FIRST_REF = 161
+// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 161]
// CK12-LABEL: implicit_maps_float_complex
void implicit_maps_float_complex (int a){
Modified: cfe/trunk/test/OpenMP/teams_distribute_firstprivate_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/teams_distribute_firstprivate_codegen.cpp?rev=314995&r1=314994&r2=314995&view=diff
==============================================================================
--- cfe/trunk/test/OpenMP/teams_distribute_firstprivate_codegen.cpp (original)
+++ cfe/trunk/test/OpenMP/teams_distribute_firstprivate_codegen.cpp Thu Oct 5 10:51:39 2017
@@ -82,6 +82,7 @@ int main() {
// LAMBDA: define{{.*}} internal{{.*}} void @[[LOFFL1]](i{{64|32}} {{%.+}}, i{{64|32}} {{%.+}})
// LAMBDA: {{%.+}} = alloca i{{[0-9]+}},
// LAMBDA: {{%.+}} = alloca i{{[0-9]+}},
+ // LAMBDA: {{%.+}} = alloca i{{[0-9]+}},
// LAMBDA: [[G_CAST:%.+]] = alloca i{{[0-9]+}},
// LAMBDA: [[SIVAR_CAST:%.+]] = alloca i{{[0-9]+}},
// LAMBDA-DAG: [[G_CAST_VAL:%.+]] = load{{.+}} [[G_CAST]],
More information about the cfe-commits
mailing list