From 2fd0cb2ae72ecb036690fcd9e0c5b9846907be16 Mon Sep 17 00:00:00 2001 From: Alexey Bataev Date: Thu, 5 Oct 2017 17:51:39 +0000 Subject: [PATCH] [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. llvm-svn: 314995 --- clang/lib/Sema/SemaOpenMP.cpp | 51 ++++++++++++++----- clang/test/OpenMP/target_map_codegen.cpp | 10 ++-- .../teams_distribute_firstprivate_codegen.cpp | 1 + 3 files changed, 43 insertions(+), 19 deletions(-) diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index 63ef808b3cc5..3938f6896844 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -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(ValueDecl *D, unsigned Level) { 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(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::ActOnOpenMPDefaultmapClause( << Value << getOpenMPClauseName(OMPC_defaultmap); return nullptr; } + DSAStack->setDefaultDMAToFromScalar(StartLoc); return new (Context) OMPDefaultmapClause(StartLoc, LParenLoc, MLoc, KindLoc, EndLoc, Kind, M); diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp index 460a02ef2dd7..72589b530f5a 100644 --- a/clang/test/OpenMP/target_map_codegen.cpp +++ b/clang/test/OpenMP/target_map_codegen.cpp @@ -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){ diff --git a/clang/test/OpenMP/teams_distribute_firstprivate_codegen.cpp b/clang/test/OpenMP/teams_distribute_firstprivate_codegen.cpp index 6529b4b6b178..1f52ebea1af9 100644 --- a/clang/test/OpenMP/teams_distribute_firstprivate_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_firstprivate_codegen.cpp @@ -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]],