[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

@@ -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<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::ActOnOpenMPDefaultmapClause(
         << Value << getOpenMPClauseName(OMPC_defaultmap);
     return nullptr;
   }
+  DSAStack->setDefaultDMAToFromScalar(StartLoc);
 
   return new (Context)
       OMPDefaultmapClause(StartLoc, LParenLoc, MLoc, KindLoc, EndLoc, Kind, M);

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){

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]],