[OpenMP] Update map clause SEMA to support OpenMP 4.5 possible list items.

[sfantao]

Summary

Extend support in the map clause SEMA for the expressions supported in the OpenMP 4.5 specification, namely member expressions.

Fix some bugs in the previous implementation of SEMA related with expressions that do not consist of single variable references.

Fix bug in parsing when the expression in the map clause do not start with an identifier: accept any expression in the map clause and check for validity in SEMA instead of just ignoring it.

[ABataev]

Please, update your workspace to latest version, many changes are in trunk already

[sfantao]

Rebase patch as requested by Alexey.

[ABataev]

LG

include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_reduction_incomplete_type : Error<
   "a reduction list item with incomplete type %0">;
 def err_omp_unexpected_clause_value : Error<
   "expected %0 in OpenMP clause '%1'">;
 def err_omp_expected_var_name_member_expr : Error<
   "expected variable name%select{| or data member of current class}0">;
 def err_omp_expected_var_name_member_expr_or_array_item : Error<
   "expected variable name%select{|, data member of current class}0, array element or array section">;
+def err_omp_expected_named_var_member_or_array_expression: Error<
+  "expected expression containing only member accesses and/or array sections based on named variables">;
+def err_omp_bit_fields_forbidden_in_map_clause : Error<
+  "bit fields cannot be used to specify storage in a map clause">;
+def err_omp_array_section_in_rightmost_expression : Error<
+  "array section can only be associated with the rightmost variable in a map clause expression">;
+def err_omp_union_type_not_allowed : Error<
+  "mapped storage cannot be derived from a union">;
+def err_omp_expected_access_to_data_field : Error<
+  "expected access to data field">;
+def err_omp_multiple_array_items_in_map_clause : Error<
+  "multiple array elements associated with the same variable are not allowed in map clauses of the same construct">;
+def err_omp_pointer_mapped_along_with_derived_section : Error<
+  "pointer cannot be mapped along with a section derived from itself">;
+def err_omp_original_storage_is_shared_and_does_not_contain : Error<
+  "original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage">;
+def err_omp_same_pointer_derreferenced : Error<
+  "same pointer derreferenced in multiple different ways in map clause expressions">;
 def note_omp_task_predetermined_firstprivate_here : Note<
   "predetermined as a firstprivate in a task construct here">;
 def err_omp_threadprivate_incomplete_type : Error<
   "threadprivate variable with incomplete type %0">;
 def err_omp_no_dsa_for_variable : Error<
   "variable %0 must have explicitly specified data sharing attributes">;
 def err_omp_wrong_dsa : Error<
   "%0 variable cannot be %1">;

lib/Parse/ParseOpenMP.cpp

OMPClause *Parser::ParseOpenMPVarListClause(OpenMPDirectiveKind DKind,
   OpenMPDependClauseKind DepKind = OMPC_DEPEND_unknown;
   // OpenMP 4.1 [2.15.3.7, linear Clause]
   //  If no modifier is specified it is assumed to be val.
   OpenMPLinearClauseKind LinearModifier = OMPC_LINEAR_val;
   OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
   OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown;
   bool MapTypeIsImplicit = false;
   bool MapTypeModifierSpecified = false;
-  bool UnexpectedId = false;
   SourceLocation DepLinMapLoc;
 
   // Parse '('.
   BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end);
   if (T.expectAndConsume(diag::err_expected_lparen_after,
                          getOpenMPClauseName(Kind)))
     return nullptr;
 
if (Tok.is(tok::identifier)) {
           MapType = OMPC_MAP_tofrom;
           MapTypeIsImplicit = true;
         }
       } else {
         MapType = OMPC_MAP_tofrom;
         MapTypeIsImplicit = true;
       }
     } else {
-      UnexpectedId = true;
+      MapType = OMPC_MAP_tofrom;
+      MapTypeIsImplicit = true;
     }
 
     if (Tok.is(tok::colon)) {
       ColonLoc = ConsumeToken();
     } else if (ColonExpected) {
       Diag(Tok, diag::warn_pragma_expected_colon) << "map type";
     }
   }
 
   SmallVector<Expr *, 5> Vars;
   bool IsComma =
       ((Kind != OMPC_reduction) && (Kind != OMPC_depend) &&
        (Kind != OMPC_map)) ||
       ((Kind == OMPC_reduction) && !InvalidReductionId) ||
-      ((Kind == OMPC_map) && (UnexpectedId || MapType != OMPC_MAP_unknown) &&
+      ((Kind == OMPC_map) && (MapType != OMPC_MAP_unknown) &&
        (!MapTypeModifierSpecified ||
         (MapTypeModifierSpecified && MapTypeModifier == OMPC_MAP_always))) ||
       ((Kind == OMPC_depend) && DepKind != OMPC_DEPEND_unknown);
   const bool MayHaveTail = (Kind == OMPC_linear || Kind == OMPC_aligned);
   while (IsComma || (Tok.isNot(tok::r_paren) && Tok.isNot(tok::colon) &&
                      Tok.isNot(tok::annot_pragma_openmp_end))) {
     ColonProtectionRAIIObject ColonRAII(*this, MayHaveTail);
     // Parse variable

lib/Sema/SemaOpenMP.cpp

struct DSAVarData {
     OpenMPClauseKind CKind;
     Expr *RefExpr;
     SourceLocation ImplicitDSALoc;
     DSAVarData()
         : DKind(OMPD_unknown), CKind(OMPC_unknown), RefExpr(nullptr),
           ImplicitDSALoc() {}
   };
 
-public:
-  struct MapInfo {
-    Expr *RefExpr;
-  };
-
 private:
+  typedef SmallVector<Expr *, 4> MapInfo;
+
   struct DSAInfo {
     OpenMPClauseKind Attributes;
     Expr *RefExpr;
   };
   typedef llvm::SmallDenseMap<ValueDecl *, DSAInfo, 64> DeclSAMapTy;
   typedef llvm::SmallDenseMap<ValueDecl *, Expr *, 64> AlignedMapTy;
   typedef llvm::DenseMap<ValueDecl *, unsigned> LoopControlVariablesMapTy;
   typedef llvm::SmallDenseMap<ValueDecl *, MapInfo, 64> MappedDeclsTy;
if (Stack.size() > 1)
       return Stack.back().InnerTeamsRegionLoc;
     return SourceLocation();
   }
 
   Scope *getCurScope() const { return Stack.back().CurScope; }
   Scope *getCurScope() { return Stack.back().CurScope; }
   SourceLocation getConstructLoc() { return Stack.back().ConstructLoc; }
 
-  MapInfo getMapInfoForVar(ValueDecl *VD) {
-    MapInfo VarMI = {0};
-    for (auto Cnt = Stack.size() - 1; Cnt > 0; --Cnt) {
-      if (Stack[Cnt].MappedDecls.count(VD)) {
-        VarMI = Stack[Cnt].MappedDecls[VD];
-        break;
-      }
+  // Do the check specified in MapInfoCheck and return true if any issue is
+  // found.
+  template <class MapInfoCheck>
+  bool checkMapInfoForVar(ValueDecl *VD, bool CurrentRegionOnly,
+                          MapInfoCheck Check) {
+    auto SI = Stack.rbegin();
+    auto SE = Stack.rend();
+
+    if (SI == SE)
+      return false;
+
+    if (CurrentRegionOnly) {
+      SE = std::next(SI);
+    } else {
+      ++SI;
+
+    for (; SI != SE; ++SI) {
+      auto MI = SI->MappedDecls.find(VD);
+      if (MI != SI->MappedDecls.end()) {
+        for (Expr *E : MI->second) {
+          if (Check(E))
+            return true;
         }
-    return VarMI;
       }
-
-  void addMapInfoForVar(ValueDecl *VD, MapInfo MI) {
-    if (Stack.size() > 1) {
-      Stack.back().MappedDecls[VD] = MI;
     }
+    return false;
   }
 
-  MapInfo IsMappedInCurrentRegion(ValueDecl *VD) {
-    assert(Stack.size() > 1 && "Target level is 0");
-    MapInfo VarMI = {0};
-    if (Stack.size() > 1 && Stack.back().MappedDecls.count(VD)) {
-      VarMI = Stack.back().MappedDecls[VD];
+  void addExprToVarMapInfo(ValueDecl *VD, Expr *E) {
+    if (Stack.size() > 1) {
+      Stack.back().MappedDecls[VD].push_back(E);
     }
-    return VarMI;
   }
 };
 bool isParallelOrTaskRegion(OpenMPDirectiveKind DKind) {
   return isOpenMPParallelDirective(DKind) || DKind == OMPD_task ||
          isOpenMPTeamsDirective(DKind) || DKind == OMPD_unknown ||
          isOpenMPTaskLoopDirective(DKind);
 }
 } // namespace
static bool CheckTypeMappable(SourceLocation SL, SourceRange SR, Sema &SemaRef,
   } else if (CXXRecordDecl *RD = dyn_cast_or_null<CXXRecordDecl>(ND)) {
     if (!RD->isInvalidDecl() &&
         !IsCXXRecordForMappable(SemaRef, SL, Stack, RD))
       return false;
   }
   return true;
 }
 
+// Return the expression of the base of the map clause or null if it cannot
+// be determined and do all the necessary checks to see if the expression is
+// valid as a standalone map clause expression.
+static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E) {
+  SourceLocation ELoc = E->getExprLoc();
+  SourceRange ERange = E->getSourceRange();
+
+  // The base of elements of list in a map clause have to be either:
+  //  - a reference to variable or field.
+  //  - a member expression.
+  //  - an array expression.
+  //
+  // E.g. if we have the expression 'r.S.Arr[:12]', we want to retrieve the
+  // reference to 'r'.
+  //
+  // If we have:
+  //
+  // struct SS {
+  //   Bla S;
+  //   foo() {
+  //     #pragma omp target map (S.Arr[:12]);
+  //   }
+  // }
+  //
+  // We want to retrieve the member expression 'this->S';
+
+  Expr *RelevantExpr = nullptr;
+
+  // Flags to help capture some memory
+
+  // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.2]
+  //  If a list item is an array section, it must specify contiguous storage.
+  //
+  // For this restriction it is sufficient that we make sure only references
+  // to variables or fields and array expressions, and that no array sections
+  // exist except in the rightmost expression. E.g. these would be invalid:
+  //
+  //   r.ArrS[3:5].Arr[6:7]
+  //
+  //   r.ArrS[3:5].x
+  //
+  // but these would be valid:
+  //   r.ArrS[3].Arr[6:7]
+  //
+  //   r.ArrS[3].x
+
+  bool IsRightMostExpression = true;
+
+  while (!RelevantExpr) {
+    auto AllowArraySection = IsRightMostExpression;
+    IsRightMostExpression = false;
+
+    E = E->IgnoreParenImpCasts();
+
+    if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
+      if (!isa<VarDecl>(CurE->getDecl()))
+        break;
+
+      RelevantExpr = CurE;
+      continue;
+    }
+
+    if (auto *CurE = dyn_cast<MemberExpr>(E)) {
+      auto *BaseE = CurE->getBase()->IgnoreParenImpCasts();
+
+      if (isa<CXXThisExpr>(BaseE))
+        // We found a base expression: this->Val.
+        RelevantExpr = CurE;
+      else
+        E = BaseE;
+
+      if (!isa<FieldDecl>(CurE->getMemberDecl())) {
+        SemaRef.Diag(ELoc, diag::err_omp_expected_access_to_data_field)
+            << CurE->getSourceRange();
+        break;
+      }
+
+      auto *FD = cast<FieldDecl>(CurE->getMemberDecl());
+
+      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.3]
+      //  A bit-field cannot appear in a map clause.
+      //
+      if (FD->isBitField()) {
+        SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_map_clause)
+            << CurE->getSourceRange();
+        break;
+      }
+
+      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
+      //  If the type of a list item is a reference to a type T then the type
+      //  will be considered to be T for all purposes of this clause.
+      QualType CurType = BaseE->getType();
+      if (CurType->isReferenceType())
+        CurType = CurType->getPointeeType();
+
+      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.2]
+      //  A list item cannot be a variable that is a member of a structure with
+      //  a union type.
+      //
+      if (auto *RT = CurType->getAs<RecordType>())
+        if (RT->isUnionType()) {
+          SemaRef.Diag(ELoc, diag::err_omp_union_type_not_allowed)
+              << CurE->getSourceRange();
+          break;
+        }
+
+      continue;
+    }
+
+    if (auto *CurE = dyn_cast<ArraySubscriptExpr>(E)) {
+      E = CurE->getBase()->IgnoreParenImpCasts();
+
+      if (!E->getType()->isAnyPointerType() && !E->getType()->isArrayType()) {
+        SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name)
+            << 0 << CurE->getSourceRange();
+        break;
+      }
+      continue;
+    }
+
+    if (auto *CurE = dyn_cast<OMPArraySectionExpr>(E)) {
+      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.7]
+      //  If a list item is an element of a structure, only the rightmost symbol
+      //  of the variable reference can be an array section.
+      //
+      if (!AllowArraySection) {
+        SemaRef.Diag(ELoc, diag::err_omp_array_section_in_rightmost_expression)
+            << CurE->getSourceRange();
+        break;
+      }
+
+      E = CurE->getBase()->IgnoreParenImpCasts();
+
+      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
+      //  If the type of a list item is a reference to a type T then the type
+      //  will be considered to be T for all purposes of this clause.
+      QualType CurType = E->getType();
+      if (CurType->isReferenceType())
+        CurType = CurType->getPointeeType();
+
+      if (!CurType->isAnyPointerType() && !CurType->isArrayType()) {
+        SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name)
+            << 0 << CurE->getSourceRange();
+        break;
+      }
+
+      continue;
+    }
+
+    // If nothing else worked, this is not a valid map clause expression.
+    SemaRef.Diag(ELoc,
+                 diag::err_omp_expected_named_var_member_or_array_expression)
+        << ERange;
+    break;
+  }
+
+  return RelevantExpr;
+}
+
+// Return true if expression E associated with value VD has conflicts with other
+// map information.
+static bool CheckMapConflicts(Sema &SemaRef, DSAStackTy *DSAS, ValueDecl *VD,
+                              Expr *E, bool CurrentRegionOnly) {
+  assert(VD && E);
+
+  // Types used to organize the components of a valid map clause.
+  typedef std::pair<Expr *, ValueDecl *> MapExpressionComponent;
+  typedef SmallVector<MapExpressionComponent, 4> MapExpressionComponents;
+
+  // Helper to extract the components in the map clause expression E and store
+  // them into MEC. This assumes that E is a valid map clause expression, i.e.
+  // it has already passed the single clause checks.
+  auto ExtractMapExpressionComponents = [](Expr *TE,
+                                           MapExpressionComponents &MEC) {
+    while (true) {
+      TE = TE->IgnoreParenImpCasts();
+
+      if (auto *CurE = dyn_cast<DeclRefExpr>(TE)) {
+        MEC.push_back(
+            MapExpressionComponent(CurE, cast<VarDecl>(CurE->getDecl())));
+        break;
+      }
+
+      if (auto *CurE = dyn_cast<MemberExpr>(TE)) {
+        auto *BaseE = CurE->getBase()->IgnoreParenImpCasts();
+
+        MEC.push_back(MapExpressionComponent(
+            CurE, cast<FieldDecl>(CurE->getMemberDecl())));
+        if (isa<CXXThisExpr>(BaseE))
+          break;
+
+        TE = BaseE;
+        continue;
+      }
+
+      if (auto *CurE = dyn_cast<ArraySubscriptExpr>(TE)) {
+        MEC.push_back(MapExpressionComponent(CurE, nullptr));
+        TE = CurE->getBase()->IgnoreParenImpCasts();
+        continue;
+      }
+
+      if (auto *CurE = dyn_cast<OMPArraySectionExpr>(TE)) {
+        MEC.push_back(MapExpressionComponent(CurE, nullptr));
+        TE = CurE->getBase()->IgnoreParenImpCasts();
+        continue;
+      }
+
+      llvm_unreachable(
+          "Expecting only valid map clause expressions at this point!");
+    }
+  };
+
+  SourceLocation ELoc = E->getExprLoc();
+  SourceRange ERange = E->getSourceRange();
+
+  // In order to easily check the conflicts we need to match each component of
+  // the expression under test with the components of the expressions that are
+  // already in the stack.
+
+  MapExpressionComponents CurComponents;
+  ExtractMapExpressionComponents(E, CurComponents);
+
+  assert(!CurComponents.empty() && "Map clause expression with no components!");
+  assert(CurComponents.back().second == VD &&
+         "Map clause expression with unexpected base!");
+
+  // Variables to help detecting enclosing problems in data environment nests.
+  bool IsEnclosedByDataEnvironmentExpr = false;
+  Expr *EnclosingExpr = nullptr;
+
+  bool FoundError =
+      DSAS->checkMapInfoForVar(VD, CurrentRegionOnly, [&](Expr *RE) -> bool {
+        MapExpressionComponents StackComponents;
+        ExtractMapExpressionComponents(RE, StackComponents);
+        assert(!StackComponents.empty() &&
+               "Map clause expression with no components!");
+        assert(StackComponents.back().second == VD &&
+               "Map clause expression with unexpected base!");
+
+        // Expressions must start from the same base. Here we detect at which
+        // point both expressions diverge from each other and see if we can
+        // detect if the memory referred to both expressions is contiguous and
+        // do not overlap.
+        auto CI = CurComponents.rbegin();
+        auto CE = CurComponents.rend();
+        auto SI = StackComponents.rbegin();
+        auto SE = StackComponents.rend();
+        for (; CI != CE && SI != SE; ++CI, ++SI) {
+
+          // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.3]
+          //  At most one list item can be an array item derived from a given
+          //  variable in map clauses of the same construct.
+          if (CurrentRegionOnly && (isa<ArraySubscriptExpr>(CI->first) ||
+                                    isa<OMPArraySectionExpr>(CI->first)) &&
+              (isa<ArraySubscriptExpr>(SI->first) ||
+               isa<OMPArraySectionExpr>(SI->first))) {
+            SemaRef.Diag(CI->first->getExprLoc(),
+                         diag::err_omp_multiple_array_items_in_map_clause)
+                << CI->first->getSourceRange();
+            ;
+            SemaRef.Diag(SI->first->getExprLoc(), diag::note_used_here)
+                << SI->first->getSourceRange();
+            return true;
+          }
+
+          // Do both expressions have the same kind?
+          if (CI->first->getStmtClass() != SI->first->getStmtClass())
+            break;
+
+          // Are we dealing with different variables/fields?
+          if (CI->second != SI->second)
+            break;
+        }
+
+        // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.4]
+        //  List items of map clauses in the same construct must not share
+        //  original storage.
+        //
+        // If the expressions are exactly the same or one is a subset of the
+        // other, it means they are sharing storage.
+        if (CI == CE && SI == SE) {
+          if (CurrentRegionOnly) {
+            SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
+            SemaRef.Diag(RE->getExprLoc(), diag::note_used_here)
+                << RE->getSourceRange();
+            return true;
+          } else {
+            // If we find the same expression in the enclosing data environment,
+            // that is legal.
+            IsEnclosedByDataEnvironmentExpr = true;
+            return false;
+          }
+        }
+
+        QualType DerivedType = std::prev(CI)->first->getType();
+        SourceLocation DerivedLoc = std::prev(CI)->first->getExprLoc();
+
+        // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
+        //  If the type of a list item is a reference to a type T then the type
+        //  will be considered to be T for all purposes of this clause.
+        if (DerivedType->isReferenceType())
+          DerivedType = DerivedType->getPointeeType();
+
+        // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.1]
+        //  A variable for which the type is pointer and an array section
+        //  derived from that variable must not appear as list items of map
+        //  clauses of the same construct.
+        //
+        // Also, cover one of the cases in:
+        // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.5]
+        //  If any part of the original storage of a list item has corresponding
+        //  storage in the device data environment, all of the original storage
+        //  must have corresponding storage in the device data environment.
+        //
+        if (DerivedType->isAnyPointerType()) {
+          if (CI == CE || SI == SE) {
+            SemaRef.Diag(
+                DerivedLoc,
+                diag::err_omp_pointer_mapped_along_with_derived_section)
+                << DerivedLoc;
+          } else {
+            assert(CI != CE && SI != SE);
+            SemaRef.Diag(DerivedLoc, diag::err_omp_same_pointer_derreferenced)
+                << DerivedLoc;
+          }
+          SemaRef.Diag(RE->getExprLoc(), diag::note_used_here)
+              << RE->getSourceRange();
+          return true;
+        }
+
+        // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.4]
+        //  List items of map clauses in the same construct must not share
+        //  original storage.
+        //
+        // An expression is a subset of the other.
+        if (CurrentRegionOnly && (CI == CE || SI == SE)) {
+          SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange;
+          SemaRef.Diag(RE->getExprLoc(), diag::note_used_here)
+              << RE->getSourceRange();
+          return true;
+        }
+
+        // The current expression uses the same base as other expression in the
+        // data environment but does not contain it completelly.
+        if (!CurrentRegionOnly && SI != SE)
+          EnclosingExpr = RE;
+
+        // The current expression is a subset of the expression in the data
+        // environment.
+        IsEnclosedByDataEnvironmentExpr |=
+            (!CurrentRegionOnly && CI != CE && SI == SE);
+
+        return false;
+      });
+
+  if (CurrentRegionOnly)
+    return FoundError;
+
+  // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.5]
+  //  If any part of the original storage of a list item has corresponding
+  //  storage in the device data environment, all of the original storage must
+  //  have corresponding storage in the device data environment.
+  // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.6]
+  //  If a list item is an element of a structure, and a different element of
+  //  the structure has a corresponding list item in the device data environment
+  //  prior to a task encountering the construct associated with the map clause,
+  //  then the list item must also have a correspnding list item in the device
+  //  data environment prior to the task encountering the construct.
+  //
+  if (EnclosingExpr && !IsEnclosedByDataEnvironmentExpr) {
+    SemaRef.Diag(ELoc,
+                 diag::err_omp_original_storage_is_shared_and_does_not_contain)
+        << ERange;
+    SemaRef.Diag(EnclosingExpr->getExprLoc(), diag::note_used_here)
+        << EnclosingExpr->getSourceRange();
+    return true;
+  }
+
+  return FoundError;
+}
+
 OMPClause *
 Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier,
                            OpenMPMapClauseKind MapType, bool IsMapTypeImplicit,
                            SourceLocation MapLoc, SourceLocation ColonLoc,
                            ArrayRef<Expr *> VarList, SourceLocation StartLoc,
                            SourceLocation LParenLoc, SourceLocation EndLoc) {
   SmallVector<Expr *, 4> Vars;
 
   for (auto &RE : VarList) {
     assert(RE && "Null expr in omp map");
     if (isa<DependentScopeDeclRefExpr>(RE)) {
       // It will be analyzed later.
       Vars.push_back(RE);
       continue;
     }
     SourceLocation ELoc = RE->getExprLoc();
 
-    // OpenMP [2.14.5, Restrictions]
-    //  A variable that is part of another variable (such as field of a
-    //  structure) but is not an array element or an array section cannot appear
-    //  in a map clause.
     auto *VE = RE->IgnoreParenLValueCasts();
 
     if (VE->isValueDependent() || VE->isTypeDependent() ||
         VE->isInstantiationDependent() ||
         VE->containsUnexpandedParameterPack()) {
-      // It will be analyzed later.
+      // We can only analyze this information once the missing information is
+      // resolved.
       Vars.push_back(RE);
       continue;
     }
 
     auto *SimpleExpr = RE->IgnoreParenCasts();
-    auto *DE = dyn_cast<DeclRefExpr>(SimpleExpr);
-    auto *ASE = dyn_cast<ArraySubscriptExpr>(SimpleExpr);
-    auto *OASE = dyn_cast<OMPArraySectionExpr>(SimpleExpr);
 
-    if (!RE->IgnoreParenImpCasts()->isLValue() ||
-        (!OASE && !ASE && !DE) ||
-        (DE && !isa<VarDecl>(DE->getDecl())) ||
-        (ASE && !ASE->getBase()->getType()->isAnyPointerType() &&
-         !ASE->getBase()->getType()->isArrayType())) {
-      Diag(ELoc, diag::err_omp_expected_var_name_member_expr_or_array_item)
-          << 0 << RE->getSourceRange();
+    if (!RE->IgnoreParenImpCasts()->isLValue()) {
+      Diag(ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
+          << RE->getSourceRange();
       continue;
     }
 
-    Decl *D = nullptr;
-    if (DE) {
-      D = DE->getDecl();
-    } else if (ASE) {
-      auto *B = ASE->getBase()->IgnoreParenCasts();
-      D = dyn_cast<DeclRefExpr>(B)->getDecl();
-    } else if (OASE) {
-      auto *B = OASE->getBase();
-      D = dyn_cast<DeclRefExpr>(B)->getDecl();
+    // Obtain the array or member expression bases if required.
+    auto *BE = CheckMapClauseExpressionBase(*this, SimpleExpr);
+    if (!BE)
+      continue;
+
+    // If the base is a reference to a variable, we rely on that variable for
+    // the following checks. If it is a 'this' expression we rely on the field.
+    ValueDecl *D = nullptr;
+    if (auto *DRE = dyn_cast<DeclRefExpr>(BE)) {
+      D = DRE->getDecl();
+    } else {
+      auto *ME = cast<MemberExpr>(BE);
+      assert(isa<CXXThisExpr>(ME->getBase()) && "Unexpected expression!");
+      D = ME->getMemberDecl();
     }
     assert(D && "Null decl on map clause.");
-    auto *VD = cast<VarDecl>(D);
 
-    // OpenMP [2.14.5, Restrictions, p.8]
-    // threadprivate variables cannot appear in a map clause.
-    if (DSAStack->isThreadPrivate(VD)) {
+    auto *VD = dyn_cast<VarDecl>(D);
+    auto *FD = dyn_cast<FieldDecl>(D);
+
+    assert((VD || FD) && "Only variables or fields are expected here!");
+
+    // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.10]
+    //  threadprivate variables cannot appear in a map clause.
+    if (VD && DSAStack->isThreadPrivate(VD)) {
       auto DVar = DSAStack->getTopDSA(VD, false);
       Diag(ELoc, diag::err_omp_threadprivate_in_map);
       ReportOriginalDSA(*this, DSAStack, VD, DVar);
       continue;
     }
 
-    // OpenMP [2.14.5, Restrictions, p.2]
-    //  At most one list item can be an array item derived from a given variable
-    //  in map clauses of the same construct.
-    // OpenMP [2.14.5, Restrictions, p.3]
-    //  List items of map clauses in the same construct must not share original
-    //  storage.
-    // OpenMP [2.14.5, Restrictions, C/C++, p.2]
-    //  A variable for which the type is pointer, reference to array, or
-    //  reference to pointer and an array section derived from that variable
-    //  must not appear as list items of map clauses of the same construct.
-    DSAStackTy::MapInfo MI = DSAStack->IsMappedInCurrentRegion(VD);
-    if (MI.RefExpr) {
-      Diag(ELoc, diag::err_omp_map_shared_storage) << ELoc;
-      Diag(MI.RefExpr->getExprLoc(), diag::note_used_here)
-          << MI.RefExpr->getSourceRange();
-      continue;
-    }
-
-    // OpenMP [2.14.5, Restrictions, C/C++, p.3,4]
-    //  A variable for which the type is pointer, reference to array, or
-    //  reference to pointer must not appear as a list item if the enclosing
-    //  device data environment already contains an array section derived from
-    //  that variable.
-    //  An array section derived from a variable for which the type is pointer,
-    //  reference to array, or reference to pointer must not appear as a list
-    //  item if the enclosing device data environment already contains that
-    //  variable.
-    QualType Type = VD->getType();
-    MI = DSAStack->getMapInfoForVar(VD);
-    if (MI.RefExpr && (isa<DeclRefExpr>(MI.RefExpr->IgnoreParenLValueCasts()) !=
-                       isa<DeclRefExpr>(VE)) &&
-        (Type->isPointerType() || Type->isReferenceType())) {
-      Diag(ELoc, diag::err_omp_map_shared_storage) << ELoc;
-      Diag(MI.RefExpr->getExprLoc(), diag::note_used_here)
-          << MI.RefExpr->getSourceRange();
-      continue;
-    }
+    // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
+    //  A list item cannot appear in both a map clause and a data-sharing
+    //  attribute clause on the same construct.
+    //
+    // TODO: Implement this check - it cannot currently be tested because of
+    // missing implementation of the other data sharing clauses in target
+    // directives.
+
+    // Check conflicts with other map clause expressions. We check the conflicts
+    // with the current construct separately from the enclosing data
+    // environment, because the restrictions are different.
+    if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr,
+                          /*CurrentRegionOnly=*/true))
+      break;
+    if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr,
+                          /*CurrentRegionOnly=*/false))
+      break;
+
+    // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
+    //  If the type of a list item is a reference to a type T then the type will
+    //  be considered to be T for all purposes of this clause.
+    QualType Type = D->getType();
+    if (Type->isReferenceType())
+      Type = Type->getPointeeType();
 
-    // OpenMP [2.14.5, Restrictions, C/C++, p.7]
+    // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9]
     //  A list item must have a mappable type.
     if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this,
                            DSAStack, Type))
       continue;
 
     // target enter data
     // OpenMP [2.10.2, Restrictions, p. 99]
     // A map-type must be specified in all map clauses and must be either
if (DKind == OMPD_target_exit_data &&
           << (IsMapTypeImplicit ? 1 : 0)
           << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
           << getOpenMPDirectiveName(DKind);
       // Proceed to add the variable in a map clause anyway, to prevent
       // further spurious messages
     }
 
     Vars.push_back(RE);
-    MI.RefExpr = RE;
-    DSAStack->addMapInfoForVar(VD, MI);
+    DSAStack->addExprToVarMapInfo(D, RE);
   }
   if (Vars.empty())
     return nullptr;
 
   return OMPMapClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars,
                               MapTypeModifier, MapType, IsMapTypeImplicit,
                               MapLoc);
 }

test/OpenMP/target_map_messages.cpp

 // RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s
 
+template <typename T, int I>
+struct SA {
+  static int ss;
+  #pragma omp threadprivate(ss) // expected-note {{defined as threadprivate or thread local}}
+  float a;
+  int b[12];
+  float *c;
+  T d;
+  float e[I];
+  T *f;
+  void func(int arg) {
+    #pragma omp target map(arg,a,d)
+    {}
+    #pragma omp target map(arg[2:2],a,d) // expected-error {{subscripted value is not an array or pointer}}
+    {}
+    #pragma omp target map(arg,a*2) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    {}
+    #pragma omp target map(arg,(c+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    {}
+    #pragma omp target map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}}
+    {}
+    #pragma omp target map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}}
+    {}
+
+    #pragma omp target map(to:ss) // expected-error {{threadprivate variables are not allowed in map clause}}
+    {}
+
+    #pragma omp target map(to:b,e)
+    {}
+    #pragma omp target map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+    {}
+    #pragma omp target map(to:b[:2],e)
+    {}
+    #pragma omp target map(to:b,e[:])
+    {}
+
+    #pragma omp target map(always, tofrom: c,f)
+    {}
+    #pragma omp target map(always, tofrom: c[1:2],f)
+    {}
+    #pragma omp target map(always, tofrom: c,f[1:2])
+    {}
+    #pragma omp target map(always, tofrom: c[:],f)   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    {}
+    #pragma omp target map(always, tofrom: c,f[:])   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    {}
+    return;
+  }
+};
+
+struct SB {
+  unsigned A;
+  unsigned B;
+  float Arr[100];
+  float *Ptr;
+  float *foo() {
+    return &Arr[0];
+  }
+};
+
+struct SC {
+  unsigned A : 2;
+  unsigned B : 3;
+  unsigned C;
+  unsigned D;
+  float Arr[100];
+  SB S;
+  SB ArrS[100];
+  SB *PtrS;
+  SB *&RPtrS;
+  float *Ptr;
+
+  SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
+};
+
+union SD {
+  unsigned A;
+  float B;
+};
+
+void SAclient(int arg) {
+  SA<int,123> s;
+  s.func(arg); // expected-note {{in instantiation of member function}}
+
+  SB *p;
+
+  SD u;
+  SC r(p),t(p);
+  #pragma omp target map(r)
+  {}
+  #pragma omp target map(r.ArrS[0].B)
+  {}
+  #pragma omp target map(r.ArrS[0].Arr[1:23])
+  {}
+  #pragma omp target map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}}
+  {}
+  #pragma omp target map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}}
+  {}
+  #pragma omp target map(r.ArrS[0].A, t.ArrS[1].A)
+  {}
+  #pragma omp target map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer derreferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+  {}
+  #pragma omp target map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer derreferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+  {}
+  #pragma omp target map(r.S.Arr[:12])
+  {}
+  #pragma omp target map(r.S.foo()[:12]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+  #pragma omp target map(r.C, r.D)
+  {}
+  #pragma omp target map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  {}
+  #pragma omp target map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  {}
+  #pragma omp target map(r.C, r.S)  // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
+  {}
+  #pragma omp target map(r, r.S)  // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  {}
+  #pragma omp target map(r.C, t.C)
+  {}
+  #pragma omp target map(r.A)   // expected-error {{bit fields cannot be used to specify storage in a map clause}}
+  {}
+  #pragma omp target map(r.Arr)
+  {}
+  #pragma omp target map(r.Arr[3:5])
+  {}
+  #pragma omp target map(r.Ptr[3:5])
+  {}
+  #pragma omp target map(r.ArrS[3:5].A)   // expected-error {{OpenMP array section is not allowed here}}
+  {}
+  #pragma omp target map(r.ArrS[3:5].Arr[6:7])   // expected-error {{OpenMP array section is not allowed here}}
+  {}
+  #pragma omp target map(r.ArrS[3].Arr[6:7])
+  {}
+  #pragma omp target map(r.S.Arr[4:5])
+  {}
+  #pragma omp target map(r.S.Ptr[4:5])
+  {}
+  #pragma omp target map(r.S.Ptr[:])  // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+  {}
+  #pragma omp target map((p+1)->A)  // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+  #pragma omp target map(u.B)  // expected-error {{mapped storage cannot be derived from a union}}
+  {}
+
+  #pragma omp target data map(to: r.C) //expected-note {{used here}}
+  {
+    #pragma omp target map(r.D)  // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
+    {}
+  }
+
+  #pragma omp target data map(to: t.Ptr) //expected-note {{used here}}
+  {
+    #pragma omp target map(t.Ptr[:23])  // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+    {}
+  }
+
+  #pragma omp target data map(to: t.C, t.D)
+  {
+  #pragma omp target data map(to: t.C)
+  {
+    #pragma omp target map(t.D)
+    {}
+  }
+  }
+
+  #pragma omp target data map(to: t)
+  {
+  #pragma omp target data map(to: t.C)
+  {
+    #pragma omp target map(t.D)
+    {}
+  }
+  }
+}
+
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 
 struct S1; // expected-note 2 {{declared here}}
T tmain(T argc) {
   T i, t[20];
   T &j = i;
   T *k = &j;
   T x;
   T y;
   T to, tofrom, always;
   const T (&l)[5] = da;
 
-
 #pragma omp target map // expected-error {{expected '(' after 'map'}}
+  {}
 #pragma omp target map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
 #pragma omp target map() // expected-error {{expected expression}}
+  {}
 #pragma omp target map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
 #pragma omp target map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
+  {}
 #pragma omp target map(to:) // expected-error {{expected expression}}
+  {}
 #pragma omp target map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  {}
 #pragma omp target map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+  {}
 #pragma omp target map(x)
   foo();
 #pragma omp target map(tofrom: t[:I])
   foo();
 #pragma omp target map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
   foo();
 #pragma omp target map(T) // expected-error {{'T' does not refer to a value}}
   foo();
-#pragma omp target map(I) // expected-error 2 {{expected variable name, array element or array section}}
+#pragma omp target map(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
   foo();
 #pragma omp target map(S2::S2s)
   foo();
 #pragma omp target map(S2::S2sc)
   foo();
 #pragma omp target map(x)
   foo();
 #pragma omp target map(to: x)
   foo();
 #pragma omp target map(to: to)
   foo();
 #pragma omp target map(to)
   foo();
 #pragma omp target map(to, x)
   foo();
-#pragma omp target map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
-#pragma omp target map(tofrom: argc > 0 ? x : y) // expected-error 2 {{expected variable name, array element or array section}} 
-#pragma omp target map(argc)
-#pragma omp target map(S1) // expected-error {{'S1' does not refer to a value}}
-#pragma omp target map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
-#pragma omp target map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
-#pragma omp target map(ca)
-#pragma omp target map(da)
-#pragma omp target map(S2::S2s)
-#pragma omp target map(S2::S2sc)
-#pragma omp target map(e, g)
-#pragma omp target map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
-#pragma omp target map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
-#pragma omp target map(k), map(k[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target data map(tofrom: argc > 0 ? x : y) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target data map(argc)
+#pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target data map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target data map(ca)
+#pragma omp target data map(da)
+#pragma omp target data map(S2::S2s)
+#pragma omp target data map(S2::S2sc)
+#pragma omp target data map(e, g)
+#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target data map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
   foo();
-#pragma omp target map(da)
+#pragma omp target data map(da)
 #pragma omp target map(da[:4])
   foo();
-#pragma omp target map(k, j, l) // expected-note 4 {{used here}}
-#pragma omp target map(k[:4]) // expected-error 2 {{variable already marked as mapped in current construct}}
-#pragma omp target map(j)
-#pragma omp target map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}}
-  foo();
-#pragma omp target map(k[:4], j, l[:5]) // expected-note 4 {{used here}}
-#pragma omp target map(k) // expected-error 2 {{variable already marked as mapped in current construct}}
-#pragma omp target map(j)
-#pragma omp target map(l) // expected-error 2 {{variable already marked as mapped in current construct}}
+#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(j)
+#pragma omp target map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+  foo();
+#pragma omp target data map(k[:4], j, l[:5]) // expected-note 4 {{used here}}
+#pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(j)
+#pragma omp target map(l) // expected-error 2 {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
   foo();
 
-#pragma omp target map(always, tofrom: x)
-#pragma omp target map(always: x) // expected-error {{missing map type}}
-#pragma omp target map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
-#pragma omp target map(always, tofrom: always, tofrom, x)
+#pragma omp target data map(always, tofrom: x)
+#pragma omp target data map(always: x) // expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
   return 0;
 }
 
 int main(int argc, char **argv) {
   const int d = 5;
   const int da[5] = { 0 };
   S4 e(4);
   S5 g(5);
   int i;
   int &j = i;
   int *k = &j;
   int x;
   int y;
   int to, tofrom, always;
   const int (&l)[5] = da;
-#pragma omp target map // expected-error {{expected '(' after 'map'}}
-#pragma omp target map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
-#pragma omp target map() // expected-error {{expected expression}}
-#pragma omp target map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
-#pragma omp target map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
-#pragma omp target map(to:) // expected-error {{expected expression}}
-#pragma omp target map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
-#pragma omp target map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target data map // expected-error {{expected '(' after 'map'}}
+#pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp target data map() // expected-error {{expected expression}}
+#pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+#pragma omp target data map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target data map(to:) // expected-error {{expected expression}}
+#pragma omp target data map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target data map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
 #pragma omp target map(x)
   foo();
 #pragma omp target map(to: x)
   foo();
 #pragma omp target map(to: to)
   foo();
 #pragma omp target map(to)
   foo();
 #pragma omp target map(to, x)
   foo();
-#pragma omp target map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
-#pragma omp target map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name, array element or array section}}
-#pragma omp target map(argc)
-#pragma omp target map(S1) // expected-error {{'S1' does not refer to a value}}
-#pragma omp target map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
-#pragma omp target map(argv[1])
-#pragma omp target map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
-#pragma omp target map(ca)
-#pragma omp target map(da)
-#pragma omp target map(S2::S2s)
-#pragma omp target map(S2::S2sc)
-#pragma omp target map(e, g)
-#pragma omp target map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
-#pragma omp target map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
-#pragma omp target map(k), map(k[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{xpected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target data map(argc)
+#pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target data map(argv[1])
+#pragma omp target data map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target data map(ca)
+#pragma omp target data map(da)
+#pragma omp target data map(S2::S2s)
+#pragma omp target data map(S2::S2sc)
+#pragma omp target data map(e, g)
+#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}}
+#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
   foo();
-#pragma omp target map(da)
+#pragma omp target data map(da)
 #pragma omp target map(da[:4])
   foo();
-#pragma omp target map(k, j, l) // expected-note 2 {{used here}}
-#pragma omp target map(k[:4]) // expected-error {{variable already marked as mapped in current construct}}
-#pragma omp target map(j)
-#pragma omp target map(l[:5]) // expected-error {{variable already marked as mapped in current construct}}
-  foo();
-#pragma omp target map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
-#pragma omp target map(k) // expected-error {{variable already marked as mapped in current construct}}
-#pragma omp target map(j)
-#pragma omp target map(l) // expected-error {{variable already marked as mapped in current construct}}
+#pragma omp target data map(k, j, l) // expected-note {{used here}}
+#pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(j)
+#pragma omp target map(l) map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  foo();
+#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(j)
+#pragma omp target map(l) // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
   foo();
 
-#pragma omp target map(always, tofrom: x)
-#pragma omp target map(always: x) // expected-error {{missing map type}}
-#pragma omp target map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
-#pragma omp target map(always, tofrom: always, tofrom, x)
+#pragma omp target data map(always, tofrom: x)
+#pragma omp target data map(always: x) // expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }