[OpenMP] Implement TR8 `present` map type modifier in Clang (1/2)

[jdenny]

Summary

This patch implements Clang front end support for the OpenMP TR8
present map type modifier. The next patch in this series implements
OpenMP runtime support.

This patch does not implement the present implicit-behavior for
defaultmap or the present motion-modifier for target update.
These can be implemented in subsequent patches.

A present map type modifier behavior that this patch does not
attempt to implement is TR8 sec. 2.22.7.1 "map Clause", p. 319,
L14-16:

If a map clause with a present map-type-modifier is present in a map
clause, then the effect of the clause is ordered before all other
map clauses that do not have the present modifier.

Compare to L10-11:

For a given construct, the effect of a map clause with the to, from,
or tofrom map-type is ordered before the effect of a map clause with
the alloc, release, or delete map-type.

I have some questions before this patch is ready for a detailed
review:

1. Is either passage quoted above relevant for any case that does not involve aliasing?

2. As far as I can tell, Clang does not implement L10-11 to handle aliasing. Is it expected that it already does? If not, then I think both passages should be implemented together later. Any objections?

3. What should the order be in the case of map(tofrom:expr1) map(present,alloc:expr2)?

4. The Clang codegen test files are difficult to maintain because of their sizes, but I tried to insert present tests into them anyway to be consistent with the existing always and close tests. I'd like to move all present codegen tests to separate test files. Any objections?

[jdenny]

Rebased.

[jdenny]

Ping. Maybe we can start with the simplest question from the review summary:

The Clang codegen test files are difficult to maintain because of their sizes, but I tried to insert present tests into them anyway to be consistent with the existing always and close tests. I'd like to move all present codegen tests to separate test files. Any objections?

[jdenny]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9582

In today's OpenMP/LLVM call, it was pointed out that this issue represents a general Clang bug for map clauses: map clauses generally shouldn't be ignored just because their variables aren't referenced in the construct.

Here, I tried to handle this issue only in the case of a present modifier because I previously didn't consider whether the existing behavior for other map types was a bug. My solution here was similar to what's done for omp target data, which doesn't have captures at all. Perhaps, this should generalized to other map types.

Another possibility might be to expand what's captured in Sema.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9582

It must be fixed in a separate patch. We definitely should not modify sema here, just the codegen. Currently, it iterates through all captures in region, also need to iterate through the explicit maps.

[jdenny]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9582

Thanks for the response. I'd be happy to generalize and extract into a new patch.

It looks like all I need to extract is the changes to generateAllInfo and emitTargetCall except I wouldn't have the PresentModifierOnly restriction. Of course, I would need to add tests. Does all that sound about right? If so, I'll work on it.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9582

Yes, something like this. Most probably, you won't need to add new tests, just modify the existing ones.

[jdenny]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9582

OK, I'm working on this. One point from the call I'm not confident about: Is this change just for globals or locals too? That is, which variables in the following examples should have map entries?

int x;
void fn() {
  int y;
  #pragma omp target map(x,y)
  {
    // no references to x and y are lexically contained
  }
}

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

9582

I think, both. We need to have references in the data environment for all explicitly mapped variables.

[jdenny]

Rebased, and extracted D83922 as discussed.

[ABataev]

This new modifier must be enabed only for OpenMP 5.1, need to add the checks for OpenMP version. Also, this change should not affect the functionality of previous versions

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7046–7047

Better to use thу next value to avoid compatibility issues with XLC.

[jdenny]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7046–7047

You mean 0x1000?

[jdenny]

Rebased onto latest D83922.

Implemented rejection of present if not OpenMP >= 5.1.

Cleaned up tests some.

[ABataev]

clang/include/clang/Basic/DiagnosticParseKinds.td

1254–1255

Better to keep original message for <= 5.0. This is what we usually do

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7046–7047

Yes

[jdenny]

clang/include/clang/Basic/DiagnosticParseKinds.td

1254–1255

What message do you mean? Which err_ diag id?

[ABataev]

clang/include/clang/Basic/DiagnosticParseKinds.td

1254–1255

For <= 5.0 the error message should be the same.

def err_omp_unknown_map_type_modifier : Error<
  "incorrect map type modifier, expected 'always', 'close', %select{or 'mapper'|'mapper', or 'presence'}0">;

In code:

Diag(..., diags::err_omp_unknown_map_type_modifier) << (LangOpts.OpenMP <= 50 ? 0 : 1);

[jdenny]

Adjusted diagnostics as requested.

Adjusted OMP_MAP_PRESENT as requested.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7046–7047

Could you add a comment that 0x800 is currently reserved for compatibility, please?

7979–7985

Why do we need this extra stuff here?

clang/lib/Parse/ParseOpenMP.cpp

3079–3080

Why do we need this?

3142

Better to use (getLangOpts().OpenMP >= 51 ? 1 : 0), the select is of integer type.

3163–3164

Same, why do we need this?

[jdenny]

clang/lib/Parse/ParseOpenMP.cpp

3079–3080

When called with OMPC_map, getOpenMPSimpleClauseType can return either an OpenMPMapModifierKind or OpenMPMapClauseKind depending on Tok. Thus, without this change, both isMapModifier and isMapType can return either of those despite the difference in their names and documentation.

I found that, when Parser::parseMapTypeModifiers ignores present as if it's not a modifier because OpenMP < 5.1, then isMapType later returns OMPC_MAP_MODIFIER_present and causes the following assert to fail in Sema::ActOnOpenMPVarListClause:

assert(0 <= ExtraModifier && ExtraModifier <= OMPC_MAP_unknown &&
           "Unexpected map modifier.");

To me, the most obvious solution is to fix isMapType and isMapModifier. Fortunately, looking in OpenMPKinds.h, the enumerator values in OpenMPMapModifierKind and OpenMPMapClauseKind are disjoint so we can tell which we have.

3163–3164

Responded on isMapModifier.

[jdenny]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7979–7985

For example:

#pragma omp target map(present, tofrom: s.x[0:3])

This generates 2 map entries:

• s: 0x1020
• s.x: 0x1000000001003

Without the above change, the 0x1000 is missing from the first entry. As a result, s is allocated, and then the presence check for s.x incorrectly passes at run time.

[jdenny]

Rebased onto a more recent master.

Added requested comment about skipping 0x800.

As requested: getLangOpts().OpenMP >= 51 -> getLangOpts().OpenMP >= 51 ? 1 : 0.

[ABataev]

clang/lib/Parse/ParseOpenMP.cpp

3079–3080

Can we have something like:

if (LangOpts.OpenMP <= 50 && Type == OMPC_MAP_MODIFIER_present)
  return OMPC_MAP_MODIFIER_unknown;

or extend getOpenMPSimpleClauseType function with the version parameter and check there is modifier is allowed and return unknown if it is not compatible with provided version?

3112–3113

Update a comment here

clang/test/OpenMP/target_data_codegen.cpp

258–260

Add a comment with interpretaion of the map flags, like OMP_TO = 0x1 | OMP_FROM=0x2 | OMP_PRESENT = 0x1000 = 0x1003

[jdenny]

clang/lib/Parse/ParseOpenMP.cpp

3079–3080

As far as I can tell, my changes general fix this bug in isMapType and isMapModifier. It makes them behave as documented. The suggestions you're making only fix them for the case of present. Why is that better?

clang/test/OpenMP/target_data_codegen.cpp

258–260

Are you asking me to add that comment to every map type encoding appearing in this patch? Or just this one?

[ABataev]

clang/lib/Parse/ParseOpenMP.cpp

3079–3080

It is too general. I think it may mask some issues in future. That's why I would suggest to tweak it for present only. Or, even better, extend getOpenMPSimpleClauseType function to handle this modifiers more correctly for each particular version rather than fix it here.

clang/test/OpenMP/target_data_codegen.cpp

258–260

It would be good to have something like this in all codegen tests you added in this patch

[jdenny]

clang/lib/Parse/ParseOpenMP.cpp

3079–3080

Or, even better, extend getOpenMPSimpleClauseType function to handle this modifiers more correctly for each particular version rather than fix it here.

One way to implement what I think you mean is to add an additional parameter to getOpenMPSimpleClauseType to request either the clause's associated type or that type's modifiers. Unless I missed a clause, this parameter would affect only map, defaultmap, and schedule clauses. For other clauses, this parameter would be ignored. Is that what you have in mind?

[jdenny]

Adjusted comments as requested.

[ABataev]

clang/lib/Parse/ParseOpenMP.cpp

3079–3080

I would also add a check for version compatibility if the modifier is supported for the given version. But anyway, I would like to take a look at what you have in mind

[jdenny]

Move present filter for OpenMP < 5.1 from parseMapTypeModifiers to getOpenMPSimpleClauseType.

Remove isMapType and isMapModifier fixes, which are no longer needed for this patch. However, add fixmes explaining how they don't always behave as documented.

clang/lib/Parse/ParseOpenMP.cpp

3079–3080

I think I misread your previous comment. I believe I've implemented your suggestion now.

[ABataev]

LG. Please, land it after the runtime part

[jdenny]

In D83061#2165063, @ABataev wrote:

LG.

Thanks for the review.

As discussed in the review summary, please consider the following. A present map type modifier behavior that this patch does not attempt to implement is TR8 sec. 2.22.7.1 "map Clause", p. 319, L14-16:

If a map clause with a present map-type-modifier is present in a map
clause, then the effect of the clause is ordered before all other
map clauses that do not have the present modifier.

Compare to L10-11:

For a given construct, the effect of a map clause with the to, from,
or tofrom map-type is ordered before the effect of a map clause with
the alloc, release, or delete map-type.

As far as I can tell, Clang does not implement L10-11. Is that correct? If not, then I think both passages should be implemented together later. Any objections?

Please, land it after the runtime part

I'll push them at the same time. The runtime patch needs to be second because its test suite depends on this patch. But there are still some details to resolve in the runtime patch as well.

[ABataev]

In D83061#2165089, @jdenny wrote:

In D83061#2165063, @ABataev wrote:

LG.

Thanks for the review.

As discussed in the review summary, please consider the following. A present map type modifier behavior that this patch does not attempt to implement is TR8 sec. 2.22.7.1 "map Clause", p. 319, L14-16:

If a map clause with a present map-type-modifier is present in a map
clause, then the effect of the clause is ordered before all other
map clauses that do not have the present modifier.

Compare to L10-11:

For a given construct, the effect of a map clause with the to, from,
or tofrom map-type is ordered before the effect of a map clause with
the alloc, release, or delete map-type.

As far as I can tell, Clang does not implement L10-11. Is that correct? If not, then I think both passages should be implemented together later. Any objections?

Looks like you're right. Yes, go ahead and implement it.

Please, land it after the runtime part

I'll push them at the same time. The runtime patch needs to be second because its test suite depends on this patch. But there are still some details to resolve in the runtime patch as well.

[jdenny]

In D83061#2165093, @ABataev wrote:

In D83061#2165089, @jdenny wrote:

In D83061#2165063, @ABataev wrote:

LG.

Thanks for the review.

As discussed in the review summary, please consider the following. A present map type modifier behavior that this patch does not attempt to implement is TR8 sec. 2.22.7.1 "map Clause", p. 319, L14-16:

If a map clause with a present map-type-modifier is present in a map
clause, then the effect of the clause is ordered before all other
map clauses that do not have the present modifier.

Compare to L10-11:

For a given construct, the effect of a map clause with the to, from,
or tofrom map-type is ordered before the effect of a map clause with
the alloc, release, or delete map-type.

As far as I can tell, Clang does not implement L10-11. Is that correct? If not, then I think both passages should be implemented together later. Any objections?

Looks like you're right. Yes, go ahead and implement it.

Are you ok for it to be a later patch after pushing these?

[ABataev]

In D83061#2165100, @jdenny wrote:

In D83061#2165093, @ABataev wrote:

In D83061#2165089, @jdenny wrote:

In D83061#2165063, @ABataev wrote:

LG.

Thanks for the review.

As discussed in the review summary, please consider the following. A present map type modifier behavior that this patch does not attempt to implement is TR8 sec. 2.22.7.1 "map Clause", p. 319, L14-16:

If a map clause with a present map-type-modifier is present in a map
clause, then the effect of the clause is ordered before all other
map clauses that do not have the present modifier.

Compare to L10-11:

For a given construct, the effect of a map clause with the to, from,
or tofrom map-type is ordered before the effect of a map clause with
the alloc, release, or delete map-type.

As far as I can tell, Clang does not implement L10-11. Is that correct? If not, then I think both passages should be implemented together later. Any objections?

Looks like you're right. Yes, go ahead and implement it.

Are you ok for it to be a later patch after pushing these?

Sure

[jdenny]

Thanks. It probably won't be soon. At the very least, I need to find the answer to question 3 from the review summary.

[jdenny]

Rebased. Due to D84182, updated the map types expected by my new tests in clang/test/OpenMP/target_data_codegen.cpp.

clang/include/clang/AST/OpenMPClause.h

class OMPMapClause final : public OMPMappableExprListClause<OMPMapClause>,
   size_t numTrailingObjects(OverloadToken<unsigned>) const {
     return getUniqueDeclarationsNum() + getTotalComponentListNum();
   }
 
 private:
   /// Map-type-modifiers for the 'map' clause.
   OpenMPMapModifierKind MapTypeModifiers[NumberOfOMPMapClauseModifiers] = {
       OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
-      OMPC_MAP_MODIFIER_unknown};
+      OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
 
   /// Location of map-type-modifiers for the 'map' clause.
   SourceLocation MapTypeModifiersLoc[NumberOfOMPMapClauseModifiers];
 
   /// Map type for the 'map' clause.
   OpenMPMapClauseKind MapType = OMPC_MAP_unknown;
 
   /// Is this an implicit map type or not.

clang/include/clang/Basic/DiagnosticParseKinds.td

def err_omp_expected_equal_in_iterator : Error<
   "expected '=' in iterator specifier">;
 def err_omp_expected_punc_after_iterator : Error<
   "expected ',' or ')' after iterator specifier">;
 def err_omp_decl_in_declare_simd_variant : Error<
   "function declaration is expected after 'declare %select{simd|variant}0' directive">;
 def err_omp_unknown_map_type : Error<
   "incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'">;
 def err_omp_unknown_map_type_modifier : Error<
-  "incorrect map type modifier, expected 'always', 'close', or 'mapper'">;
+  "incorrect map type modifier, expected 'always', 'close', "
+  "%select{or 'mapper'|'mapper', or 'present'}0">;
 def err_omp_map_type_missing : Error<
   "missing map type">;
 def err_omp_map_type_modifier_missing : Error<
   "missing map type modifier">;
 def err_omp_declare_simd_inbranch_notinbranch : Error<
   "unexpected '%0' clause, '%1' is specified already">;
 def err_expected_end_declare_target_or_variant : Error<
   "expected '#pragma omp end declare %select{target|variant}0'">;
 def err_expected_begin_declare_variant

clang/include/clang/Basic/OpenMPKinds.h

 
 /// OpenMP modifiers for 'reduction' clause.
 enum OpenMPReductionClauseModifier {
 #define OPENMP_REDUCTION_MODIFIER(Name) OMPC_REDUCTION_##Name,
 #include "clang/Basic/OpenMPKinds.def"
   OMPC_REDUCTION_unknown,
 };
 
-unsigned getOpenMPSimpleClauseType(OpenMPClauseKind Kind, llvm::StringRef Str);
+unsigned getOpenMPSimpleClauseType(OpenMPClauseKind Kind, llvm::StringRef Str,
+                                   unsigned OpenMPVersion);
 const char *getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, unsigned Type);
 
 /// Checks if the specified directive is a directive with an associated
 /// loop construct.
 /// \param DKind Specified directive.
 /// \return true - the directive is a loop-associated directive like 'omp simd'
 /// or 'omp for' directive, otherwise - false.
 bool isOpenMPLoopDirective(OpenMPDirectiveKind DKind);

clang/include/clang/Basic/OpenMPKinds.def

 OPENMP_MAP_KIND(tofrom)
 OPENMP_MAP_KIND(delete)
 OPENMP_MAP_KIND(release)
 
 // Map-type-modifiers for 'map' clause.
 OPENMP_MAP_MODIFIER_KIND(always)
 OPENMP_MAP_MODIFIER_KIND(close)
 OPENMP_MAP_MODIFIER_KIND(mapper)
+OPENMP_MAP_MODIFIER_KIND(present)
 
 // Modifiers for 'to' clause.
 OPENMP_TO_MODIFIER_KIND(mapper)
 
 // Modifiers for 'from' clause.
 OPENMP_FROM_MODIFIER_KIND(mapper)
 
 // Static attributes for 'dist_schedule' clause.

clang/lib/Basic/OpenMPKinds.cpp

 #include "llvm/ADT/StringRef.h"
 #include "llvm/ADT/StringSwitch.h"
 #include "llvm/Support/ErrorHandling.h"
 #include <cassert>
 
 using namespace clang;
 using namespace llvm::omp;
 
-unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind,
-                                          StringRef Str) {
+unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
+                                          unsigned OpenMPVersion) {
   switch (Kind) {
   case OMPC_default:
     return llvm::StringSwitch<unsigned>(Str)
 #define OMP_DEFAULT_KIND(Enum, Name) .Case(Name, unsigned(Enum))
 #include "llvm/Frontend/OpenMP/OMPKinds.def"
         .Default(unsigned(llvm::omp::OMP_DEFAULT_unknown));
   case OMPC_proc_bind:
     return llvm::StringSwitch<unsigned>(Str)

 #define OPENMP_DEPEND_KIND(Name) .Case(#Name, OMPC_DEPEND_##Name)
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_DEPEND_unknown);
   case OMPC_linear:
     return llvm::StringSwitch<OpenMPLinearClauseKind>(Str)
 #define OPENMP_LINEAR_KIND(Name) .Case(#Name, OMPC_LINEAR_##Name)
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_LINEAR_unknown);
-  case OMPC_map:
-    return llvm::StringSwitch<unsigned>(Str)
+  case OMPC_map: {
+    unsigned Type = llvm::StringSwitch<unsigned>(Str)
 #define OPENMP_MAP_KIND(Name)                                                  \
   .Case(#Name, static_cast<unsigned>(OMPC_MAP_##Name))
 #define OPENMP_MAP_MODIFIER_KIND(Name)                                         \
   .Case(#Name, static_cast<unsigned>(OMPC_MAP_MODIFIER_##Name))
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_MAP_unknown);
+    if (OpenMPVersion < 51 && Type == OMPC_MAP_MODIFIER_present)
+      return OMPC_MAP_MODIFIER_unknown;
+    return Type;
+  }
   case OMPC_to:
     return llvm::StringSwitch<unsigned>(Str)
 #define OPENMP_TO_MODIFIER_KIND(Name)                                          \
   .Case(#Name, static_cast<unsigned>(OMPC_TO_MODIFIER_##Name))
 #include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_TO_MODIFIER_unknown);
   case OMPC_from:
     return llvm::StringSwitch<unsigned>(Str)

clang/lib/CodeGen/CGOpenMPRuntime.cpp

enum OpenMPOffloadMappingFlags : uint64_t {
     OMP_MAP_PRIVATE = 0x80,
     /// Pass the element to the device by value.
     OMP_MAP_LITERAL = 0x100,
     /// Implicit map
     OMP_MAP_IMPLICIT = 0x200,
     /// Close is a hint to the runtime to allocate memory close to
     /// the target device.
     OMP_MAP_CLOSE = 0x400,
+    /// 0x800 is reserved for compatibility with XLC.
+    /// Produce a runtime error if the data is not already allocated.
+    OMP_MAP_PRESENT = 0x1000,
     /// The 16 MSBs of the flags indicate whether the entry is member of some
     /// struct/class.
     OMP_MAP_MEMBER_OF = 0xffff000000000000,
     LLVM_MARK_AS_BITMASK_ENUM(/* LargestFlag = */ OMP_MAP_MEMBER_OF),
   };
 
   /// Get the offset of the OMP_MAP_MEMBER_OF field.
   static unsigned getFlagMemberOffset() {
OpenMPOffloadMappingFlags getMapTypeBits(
     if (AddIsTargetParamFlag)
       Bits |= OMP_MAP_TARGET_PARAM;
     if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_always)
         != MapModifiers.end())
       Bits |= OMP_MAP_ALWAYS;
     if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
         != MapModifiers.end())
       Bits |= OMP_MAP_CLOSE;
+    if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present)
+        != MapModifiers.end())
+      Bits |= OMP_MAP_PRESENT;
     return Bits;
   }
 
   /// Return true if the provided expression is a final array section. A
   /// final array section, is one whose length can't be proved to be one.
   bool isFinalArraySectionExpression(const Expr *E) const {
     const auto *OASE = dyn_cast<OMPArraySectionExpr>(E);
 
void emitCombinedEntry(MapCombinedInfoTy &CombinedInfo,
     llvm::Value *CLAddr = CGF.Builder.CreatePointerCast(LB, CGF.VoidPtrTy);
     llvm::Value *CHAddr = CGF.Builder.CreatePointerCast(HAddr, CGF.VoidPtrTy);
     llvm::Value *Diff = CGF.Builder.CreatePtrDiff(CHAddr, CLAddr);
     llvm::Value *Size = CGF.Builder.CreateIntCast(Diff, CGF.Int64Ty,
                                                   /*isSigned=*/false);
     CombinedInfo.Sizes.push_back(Size);
     // Map type is always TARGET_PARAM
     CombinedInfo.Types.push_back(OMP_MAP_TARGET_PARAM);
+    // If any element has the present modifier, then make sure the runtime
+    // doesn't attempt to allocate the struct.
+    if (CurTypes.end() !=
+        llvm::find_if(CurTypes, [](OpenMPOffloadMappingFlags Type) {
+          return Type & OMP_MAP_PRESENT;
+        }))
+      CombinedInfo.Types.back() |= OMP_MAP_PRESENT;
     // Remove TARGET_PARAM flag from the first element
     (*CurTypes.begin()) &= ~OMP_MAP_TARGET_PARAM;
 
     // All other current entries will be MEMBER_OF the combined entry
     // (except for PTR_AND_OBJ entries which do not have a placeholder value
     // 0xFFFF in the MEMBER_OF field).
     OpenMPOffloadMappingFlags MemberOfFlag =
         getMemberOfFlag(CombinedInfo.BasePointers.size() - 1);
for (CapturedStmt::const_capture_iterator CI = CS.capture_begin(),
       CombinedInfo.append(CurInfo);
     }
     // Adjust MEMBER_OF flags for the lambdas captures.
     MEHandler.adjustMemberOfForLambdaCaptures(
         LambdaPointers, CombinedInfo.BasePointers, CombinedInfo.Pointers,
         CombinedInfo.Types);
     // Map any list items in a map clause that were not captures because they
     // weren't referenced within the construct.
     MEHandler.generateAllInfo(CombinedInfo, MappedVarSet);
 
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
     emitOffloadingArrays(CGF, CombinedInfo, Info);
     emitOffloadingArraysArgument(CGF, Info.BasePointersArray,
                                  Info.PointersArray, Info.SizesArray,
                                  Info.MapTypesArray, Info.MappersArray, Info);
     InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;

clang/lib/Parse/ParseOpenMP.cpp

parseOpenMPSimpleClause(Parser &P, OpenMPClauseKind Kind) {
   SourceLocation LOpen = P.ConsumeToken();
   // Parse '('.
   BalancedDelimiterTracker T(P, tok::l_paren, tok::annot_pragma_openmp_end);
   if (T.expectAndConsume(diag::err_expected_lparen_after,
                          getOpenMPClauseName(Kind).data()))
     return llvm::None;
 
   unsigned Type = getOpenMPSimpleClauseType(
-      Kind, Tok.isAnnotation() ? "" : P.getPreprocessor().getSpelling(Tok));
+      Kind, Tok.isAnnotation() ? "" : P.getPreprocessor().getSpelling(Tok),
+      P.getLangOpts().OpenMP);
   SourceLocation TypeLoc = Tok.getLocation();
   if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
       Tok.isNot(tok::annot_pragma_openmp_end))
     P.ConsumeAnyToken();
 
   // Parse ')'.
   SourceLocation RLoc = Tok.getLocation();
   if (!T.consumeClose())
OMPClause *Parser::ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
   if (Kind == OMPC_schedule) {
     enum { Modifier1, Modifier2, ScheduleKind, NumberOfElements };
     Arg.resize(NumberOfElements);
     KLoc.resize(NumberOfElements);
     Arg[Modifier1] = OMPC_SCHEDULE_MODIFIER_unknown;
     Arg[Modifier2] = OMPC_SCHEDULE_MODIFIER_unknown;
     Arg[ScheduleKind] = OMPC_SCHEDULE_unknown;
     unsigned KindModifier = getOpenMPSimpleClauseType(
-        Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok));
+        Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
+        getLangOpts().OpenMP);
     if (KindModifier > OMPC_SCHEDULE_unknown) {
       // Parse 'modifier'
       Arg[Modifier1] = KindModifier;
       KLoc[Modifier1] = Tok.getLocation();
       if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
           Tok.isNot(tok::annot_pragma_openmp_end))
         ConsumeAnyToken();
       if (Tok.is(tok::comma)) {
         // Parse ',' 'modifier'
         ConsumeAnyToken();
         KindModifier = getOpenMPSimpleClauseType(
-            Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok));
+            Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
+            getLangOpts().OpenMP);
         Arg[Modifier2] = KindModifier > OMPC_SCHEDULE_unknown
                              ? KindModifier
                              : (unsigned)OMPC_SCHEDULE_unknown;
         KLoc[Modifier2] = Tok.getLocation();
         if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
             Tok.isNot(tok::annot_pragma_openmp_end))
           ConsumeAnyToken();
       }
       // Parse ':'
       if (Tok.is(tok::colon))
         ConsumeAnyToken();
       else
         Diag(Tok, diag::warn_pragma_expected_colon) << "schedule modifier";
       KindModifier = getOpenMPSimpleClauseType(
-          Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok));
+          Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
+          getLangOpts().OpenMP);
     }
     Arg[ScheduleKind] = KindModifier;
     KLoc[ScheduleKind] = Tok.getLocation();
     if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
         Tok.isNot(tok::annot_pragma_openmp_end))
       ConsumeAnyToken();
     if ((Arg[ScheduleKind] == OMPC_SCHEDULE_static ||
          Arg[ScheduleKind] == OMPC_SCHEDULE_dynamic ||
          Arg[ScheduleKind] == OMPC_SCHEDULE_guided) &&
         Tok.is(tok::comma))
       DelimLoc = ConsumeAnyToken();
   } else if (Kind == OMPC_dist_schedule) {
     Arg.push_back(getOpenMPSimpleClauseType(
-        Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok)));
+        Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
+        getLangOpts().OpenMP));
     KLoc.push_back(Tok.getLocation());
     if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
         Tok.isNot(tok::annot_pragma_openmp_end))
       ConsumeAnyToken();
     if (Arg.back() == OMPC_DIST_SCHEDULE_static && Tok.is(tok::comma))
       DelimLoc = ConsumeAnyToken();
   } else if (Kind == OMPC_defaultmap) {
     // Get a defaultmap modifier
     unsigned Modifier = getOpenMPSimpleClauseType(
-        Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok));
+        Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
+        getLangOpts().OpenMP);
     // Set defaultmap modifier to unknown if it is either scalar, aggregate, or
     // pointer
     if (Modifier < OMPC_DEFAULTMAP_MODIFIER_unknown)
       Modifier = OMPC_DEFAULTMAP_MODIFIER_unknown;
     Arg.push_back(Modifier);
     KLoc.push_back(Tok.getLocation());
     if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
         Tok.isNot(tok::annot_pragma_openmp_end))
       ConsumeAnyToken();
     // Parse ':'
     if (Tok.is(tok::colon) || getLangOpts().OpenMP < 50) {
       if (Tok.is(tok::colon))
         ConsumeAnyToken();
       else if (Arg.back() != OMPC_DEFAULTMAP_MODIFIER_unknown)
         Diag(Tok, diag::warn_pragma_expected_colon) << "defaultmap modifier";
       // Get a defaultmap kind
       Arg.push_back(getOpenMPSimpleClauseType(
-          Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok)));
+          Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
+          getLangOpts().OpenMP));
       KLoc.push_back(Tok.getLocation());
       if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) &&
           Tok.isNot(tok::annot_pragma_openmp_end))
         ConsumeAnyToken();
     } else {
       Arg.push_back(OMPC_DEFAULTMAP_unknown);
       KLoc.push_back(SourceLocation());
     }
   } else if (Kind == OMPC_device) {
     // Only target executable directives support extended device construct.
     if (isOpenMPTargetExecutionDirective(DKind) && getLangOpts().OpenMP >= 50 &&
         NextToken().is(tok::colon)) {
       // Parse optional <device modifier> ':'
       Arg.push_back(getOpenMPSimpleClauseType(
-          Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok)));
+          Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok),
+          getLangOpts().OpenMP));
       KLoc.push_back(Tok.getLocation());
       ConsumeAnyToken();
       // Parse ':'
       ConsumeAnyToken();
     } else {
       Arg.push_back(OMPC_DEVICE_unknown);
       KLoc.emplace_back();
     }
return P.ParseUnqualifiedId(
       ReductionIdScopeSpec, /*ObjectType=*/nullptr,
       /*ObjectHadErrors=*/false, /*EnteringContext*/ false,
       /*AllowDestructorName*/ false,
       /*AllowConstructorName*/ false,
       /*AllowDeductionGuide*/ false, nullptr, ReductionId);
 }
 
 /// Checks if the token is a valid map-type-modifier.
+/// FIXME: It will return an OpenMPMapClauseKind if that's what it parses.
 static OpenMPMapModifierKind isMapModifier(Parser &P) {
   Token Tok = P.getCurToken();
   if (!Tok.is(tok::identifier))
     return OMPC_MAP_MODIFIER_unknown;
 
   Preprocessor &PP = P.getPreprocessor();
-  OpenMPMapModifierKind TypeModifier = static_cast<OpenMPMapModifierKind>(
-      getOpenMPSimpleClauseType(OMPC_map, PP.getSpelling(Tok)));
+  OpenMPMapModifierKind TypeModifier =
+      static_cast<OpenMPMapModifierKind>(getOpenMPSimpleClauseType(
+          OMPC_map, PP.getSpelling(Tok), P.getLangOpts().OpenMP));
   return TypeModifier;
 }
 
 /// Parse the mapper modifier in map, to, and from clauses.
 bool Parser::parseMapperModifier(OpenMPVarListDataTy &Data) {
   // Parse '('.
   BalancedDelimiterTracker T(*this, tok::l_paren, tok::colon);
   if (T.expectAndConsume(diag::err_expected_lparen_after, "mapper")) {
     SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end,
               StopBeforeMatch);
     return true;
Data.ReductionOrMapperId = DeclarationNameInfo(
       DeclNames.getIdentifier(Tok.getIdentifierInfo()), Tok.getLocation());
   ConsumeToken();
   // Parse ')'.
   return T.consumeClose();
 }
 
 /// Parse map-type-modifiers in map clause.
 /// map([ [map-type-modifier[,] [map-type-modifier[,] ...] map-type : ] list)
-/// where, map-type-modifier ::= always | close | mapper(mapper-identifier)
+/// where, map-type-modifier ::= always | close | mapper(mapper-identifier) |
+/// present
 bool Parser::parseMapTypeModifiers(OpenMPVarListDataTy &Data) {
   while (getCurToken().isNot(tok::colon)) {
     OpenMPMapModifierKind TypeModifier = isMapModifier(*this);
     if (TypeModifier == OMPC_MAP_MODIFIER_always ||
-        TypeModifier == OMPC_MAP_MODIFIER_close) {
+        TypeModifier == OMPC_MAP_MODIFIER_close ||
+        TypeModifier == OMPC_MAP_MODIFIER_present) {
       Data.MapTypeModifiers.push_back(TypeModifier);
       Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
       ConsumeToken();
     } else if (TypeModifier == OMPC_MAP_MODIFIER_mapper) {
       Data.MapTypeModifiers.push_back(TypeModifier);
       Data.MapTypeModifiersLoc.push_back(Tok.getLocation());
       ConsumeToken();
       if (parseMapperModifier(Data))
         return true;
     } else {
       // For the case of unknown map-type-modifier or a map-type.
       // Map-type is followed by a colon; the function returns when it
       // encounters a token followed by a colon.
       if (Tok.is(tok::comma)) {
         Diag(Tok, diag::err_omp_map_type_modifier_missing);
         ConsumeToken();
         continue;
       }
       // Potential map-type token as it is followed by a colon.
       if (PP.LookAhead(0).is(tok::colon))
         return false;
-      Diag(Tok, diag::err_omp_unknown_map_type_modifier);
+      Diag(Tok, diag::err_omp_unknown_map_type_modifier)
+          << (getLangOpts().OpenMP >= 51 ? 1 : 0);
       ConsumeToken();
     }
     if (getCurToken().is(tok::comma))
       ConsumeToken();
   }
   return false;
 }
 
 /// Checks if the token is a valid map-type.
+/// FIXME: It will return an OpenMPMapModifierKind if that's what it parses.
 static OpenMPMapClauseKind isMapType(Parser &P) {
   Token Tok = P.getCurToken();
   // The map-type token can be either an identifier or the C++ delete keyword.
   if (!Tok.isOneOf(tok::identifier, tok::kw_delete))
     return OMPC_MAP_unknown;
   Preprocessor &PP = P.getPreprocessor();
-  OpenMPMapClauseKind MapType = static_cast<OpenMPMapClauseKind>(
-      getOpenMPSimpleClauseType(OMPC_map, PP.getSpelling(Tok)));
+  OpenMPMapClauseKind MapType =
+      static_cast<OpenMPMapClauseKind>(getOpenMPSimpleClauseType(
+          OMPC_map, PP.getSpelling(Tok), P.getLangOpts().OpenMP));
   return MapType;
 }
 
 /// Parse map-type in map clause.
 /// map([ [map-type-modifier[,] [map-type-modifier[,] ...] map-type : ] list)
 /// where, map-type ::= to | from | tofrom | alloc | release | delete
 static void parseMapType(Parser &P, Parser::OpenMPVarListDataTy &Data) {
   Token Tok = P.getCurToken();
   if (Tok.is(tok::colon)) {
     P.Diag(Tok, diag::err_omp_map_type_missing);
     return;
bool Parser::ParseOpenMPVarList(OpenMPDirectiveKind DKind,
   // Handle reduction-identifier for reduction clause.
   if (Kind == OMPC_reduction || Kind == OMPC_task_reduction ||
       Kind == OMPC_in_reduction) {
     Data.ExtraModifier = OMPC_REDUCTION_unknown;
     if (Kind == OMPC_reduction && getLangOpts().OpenMP >= 50 &&
         (Tok.is(tok::identifier) || Tok.is(tok::kw_default)) &&
         NextToken().is(tok::comma)) {
       // Parse optional reduction modifier.
-      Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok));
+      Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok),
+                                                     getLangOpts().OpenMP);
       Data.ExtraModifierLoc = Tok.getLocation();
       ConsumeToken();
       assert(Tok.is(tok::comma) && "Expected comma.");
       (void)ConsumeToken();
     }
     ColonProtectionRAIIObject ColonRAII(*this);
     if (getLangOpts().CPlusPlus)
       ParseOptionalCXXScopeSpecifier(Data.ReductionOrMapperIdScopeSpec,
if (getLangOpts().OpenMP >= 50) {
         Data.DepModOrTailExpr = IteratorRes.get();
         // Parse ','
         ExpectAndConsume(tok::comma);
       }
     }
     // Handle dependency type for depend clause.
     ColonProtectionRAIIObject ColonRAII(*this);
     Data.ExtraModifier = getOpenMPSimpleClauseType(
-        Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : "");
+        Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : "",
+        getLangOpts().OpenMP);
     Data.ExtraModifierLoc = Tok.getLocation();
     if (Data.ExtraModifier == OMPC_DEPEND_unknown) {
       SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end,
                 StopBeforeMatch);
     } else {
       ConsumeToken();
       // Special processing for depend(source) clause.
       if (DKind == OMPD_ordered && Data.ExtraModifier == OMPC_DEPEND_source) {
         // Parse ')'.
         T.consumeClose();
         return false;
       }
     }
     if (Tok.is(tok::colon)) {
       Data.ColonLoc = ConsumeToken();
     } else {
       Diag(Tok, DKind == OMPD_ordered ? diag::warn_pragma_expected_colon_r_paren
                                       : diag::warn_pragma_expected_colon)
           << "dependency type";
     }
   } else if (Kind == OMPC_linear) {
     // Try to parse modifier if any.
     Data.ExtraModifier = OMPC_LINEAR_val;
     if (Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::l_paren)) {
-      Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok));
+      Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok),
+                                                     getLangOpts().OpenMP);
       Data.ExtraModifierLoc = ConsumeToken();
       LinearT.consumeOpen();
       NeedRParenForLinear = true;
     }
   } else if (Kind == OMPC_lastprivate) {
     // Try to parse modifier if any.
     Data.ExtraModifier = OMPC_LASTPRIVATE_unknown;
     // Conditional modifier allowed only in OpenMP 5.0 and not supported in
     // distribute and taskloop based directives.
     if ((getLangOpts().OpenMP >= 50 && !isOpenMPDistributeDirective(DKind) &&
          !isOpenMPTaskLoopDirective(DKind)) &&
         Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::colon)) {
-      Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok));
+      Data.ExtraModifier = getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok),
+                                                     getLangOpts().OpenMP);
       Data.ExtraModifierLoc = Tok.getLocation();
       ConsumeToken();
       assert(Tok.is(tok::colon) && "Expected colon.");
       Data.ColonLoc = ConsumeToken();
     }
   } else if (Kind == OMPC_map) {
     // Handle map type for map clause.
     ColonProtectionRAIIObject ColonRAII(*this);
if (Kind == OMPC_reduction || Kind == OMPC_task_reduction ||
     }
 
     if (Tok.is(tok::colon))
       Data.ColonLoc = ConsumeToken();
   } else if (Kind == OMPC_to || Kind == OMPC_from) {
     if (Tok.is(tok::identifier)) {
       bool IsMapperModifier = false;
       if (Kind == OMPC_to) {
-        auto Modifier = static_cast<OpenMPToModifierKind>(
-            getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok)));
+        auto Modifier =
+            static_cast<OpenMPToModifierKind>(getOpenMPSimpleClauseType(
+                Kind, PP.getSpelling(Tok), getLangOpts().OpenMP));
         if (Modifier == OMPC_TO_MODIFIER_mapper)
           IsMapperModifier = true;
       } else {
-        auto Modifier = static_cast<OpenMPFromModifierKind>(
-            getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok)));
+        auto Modifier =
+            static_cast<OpenMPFromModifierKind>(getOpenMPSimpleClauseType(
+                Kind, PP.getSpelling(Tok), getLangOpts().OpenMP));
         if (Modifier == OMPC_FROM_MODIFIER_mapper)
           IsMapperModifier = true;
       }
       if (IsMapperModifier) {
         // Parse the mapper modifier.
         ConsumeToken();
         IsInvalidMapperModifier = parseMapperModifier(Data);
         if (Tok.isNot(tok::colon)) {

clang/lib/Sema/SemaOpenMP.cpp

 
 OMPClause *Sema::ActOnOpenMPMapClause(
     ArrayRef<OpenMPMapModifierKind> MapTypeModifiers,
     ArrayRef<SourceLocation> MapTypeModifiersLoc,
     CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId,
     OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, SourceLocation MapLoc,
     SourceLocation ColonLoc, ArrayRef<Expr *> VarList,
     const OMPVarListLocTy &Locs, ArrayRef<Expr *> UnresolvedMappers) {
-  OpenMPMapModifierKind Modifiers[] = {OMPC_MAP_MODIFIER_unknown,
-                                       OMPC_MAP_MODIFIER_unknown,
-                                       OMPC_MAP_MODIFIER_unknown};
+  OpenMPMapModifierKind Modifiers[] = {
+      OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown,
+      OMPC_MAP_MODIFIER_unknown, OMPC_MAP_MODIFIER_unknown};
   SourceLocation ModifiersLoc[NumberOfOMPMapClauseModifiers];
 
   // Process map-type-modifiers, flag errors for duplicate modifiers.
   unsigned Count = 0;
   for (unsigned I = 0, E = MapTypeModifiers.size(); I < E; ++I) {
     if (MapTypeModifiers[I] != OMPC_MAP_MODIFIER_unknown &&
         llvm::find(Modifiers, MapTypeModifiers[I]) != std::end(Modifiers)) {
       Diag(MapTypeModifiersLoc[I], diag::err_omp_duplicate_map_type_modifier);

clang/test/OpenMP/target_data_ast_print.cpp

 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
 
 // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
+
+// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s
 // expected-no-diagnostics
 
 #ifndef HEADER
 #define HEADER
 
 void foo() {}
 
 template <typename T, int C>
#pragma omp target data map(to: c) map(from: d)
   foo();
 
 #pragma omp target data map(always,alloc: e)
   foo();
 
 #pragma omp target data map(close,alloc: e)
   foo();
 
+#ifdef OMP51
+#pragma omp target data map(present,alloc: e)
+  foo();
+#endif
+
 // nesting a target region
 #pragma omp target data map(e)
 {
   #pragma omp target map(always, alloc: e)
     foo();
   #pragma omp target map(close, alloc: e)
     foo();
+#ifdef OMP51
+  #pragma omp target map(present, alloc: e)
+    foo();
+#endif
 }
 
   return 0;
 }
 
 // CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
 // CHECK-NEXT: T i, j, b, c, d, e, x[20];
 // CHECK-NEXT: #pragma omp target data map(to: c){{$}}

 // CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(to: c) map(from: d)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(always,alloc: e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(close,alloc: e)
 // CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
+// OMP51-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: e)
 // CHECK-NEXT: {
 // CHECK-NEXT: #pragma omp target map(always,alloc: e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target map(close,alloc: e)
 // CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target map(present,alloc: e)
+// OMP51-NEXT: foo();
 // CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
 // CHECK-NEXT: int i, j, b, c, d, e, x[20];
 // CHECK-NEXT: #pragma omp target data map(to: c)
 // CHECK-NEXT: i = argc;
 // CHECK-NEXT: #pragma omp target data map(to: c) if(target data: j > 0)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(to: c) if(b)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: c)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(to: c) map(from: d)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(always,alloc: e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(close,alloc: e)
 // CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
+// OMP51-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: e)
 // CHECK-NEXT: {
 // CHECK-NEXT: #pragma omp target map(always,alloc: e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target map(close,alloc: e)
 // CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target map(present,alloc: e)
+// OMP51-NEXT: foo();
 // CHECK: template<> char tmain<char, 1>(char argc, char *argv) {
 // CHECK-NEXT: char i, j, b, c, d, e, x[20];
 // CHECK-NEXT: #pragma omp target data map(to: c)
 // CHECK-NEXT: i = argc;
 // CHECK-NEXT: #pragma omp target data map(to: c) if(target data: j > 0)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(to: c) if(b)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: c)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(to: c) map(from: d)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(always,alloc: e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(close,alloc: e)
 // CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
+// OMP51-NEXT: foo();
 // CHECK-NEXT: #pragma omp target data map(tofrom: e)
 // CHECK-NEXT: {
 // CHECK-NEXT: #pragma omp target map(always,alloc: e)
 // CHECK-NEXT: foo();
 // CHECK-NEXT: #pragma omp target map(close,alloc: e)
 // CHECK-NEXT: foo();
+// OMP51-NEXT: #pragma omp target map(present,alloc: e)
+// OMP51-NEXT: foo();
 
 int main (int argc, char **argv) {
   int b = argc, c, d, e, f, g, x[20];
   static int a;
 // CHECK: static int a;
 
 #pragma omp target data map(to: ([argc][3][a])argv)
   // CHECK: #pragma omp target data map(to: ([argc][3][a])argv)
// CHECK-NEXT: #pragma omp target data map(always,alloc: e)
   foo();
 // CHECK-NEXT: foo();
 
 #pragma omp target data map(close,alloc: e)
 // CHECK-NEXT: #pragma omp target data map(close,alloc: e)
   foo();
 // CHECK-NEXT: foo();
 
+// OMP51-NEXT: #pragma omp target data map(present,alloc: e)
+// OMP51-NEXT: foo();
+#ifdef OMP51
+#pragma omp target data map(present,alloc: e)
+  foo();
+#endif
+
 // nesting a target region
 #pragma omp target data map(e)
 // CHECK-NEXT: #pragma omp target data map(tofrom: e)
 {
 // CHECK-NEXT: {
   #pragma omp target map(always, alloc: e)
 // CHECK-NEXT: #pragma omp target map(always,alloc: e)
     foo();
 // CHECK-NEXT: foo();
 #pragma omp target map(close, alloc: e)
 // CHECK-NEXT: #pragma omp target map(close,alloc: e)
   foo();
+// CHECK-NEXT: foo();
+#pragma omp target map(always, alloc: e)
+// CHECK-NEXT: #pragma omp target map(always,alloc: e)
+  foo();
 }
 
   return tmain<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]);
 }
 
 #endif

clang/test/OpenMP/target_data_codegen.cpp

// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
 
   // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
   // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
   // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
   // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
   #pragma omp target data map(always close, to: lb)
   {++arg;}
+
+}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-64
+// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-64
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-32
+// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-32
+
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK1A
+
+// CK1A: [[ST:%.+]] = type { i32, double* }
+template <typename T>
+struct ST {
+  T a;
+  double *b;
+};
+
+ST<int> gb;
+double gc[100];
+
+// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
+// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+
+// PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x1425
+// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+
+// CK1A-LABEL: _Z3fooi
+void foo(int arg) {
+  int la;
+  float lb[arg];
+
+  // Region 00
+  // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+  // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
+  // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+  // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+  // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+  // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+  // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+
+  // CK1A-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
+  // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
+  #pragma omp target data map(present, to: lb)
+  {++arg;}
+
+  // Region 01
+  // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+  // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
+  // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+  // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+  // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+  // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+  // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+
+  // CK1A-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P]]
+  // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S]]
+  #pragma omp target data map(always close present, to: lb)
+  {++arg;}
+
 }
 #endif
 ///==========================================================================///
 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-64
 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s

 void test_device_ptr_addr(int arg) {
   int *p;
   // CK7: add nsw i32
   // CK7: add nsw i32
   #pragma omp target data use_device_ptr(p) use_device_addr(arg)
   { ++arg, ++(*p); }
 }
 #endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK8 --check-prefix CK8-64
+// RUN: %clang_cc1 -DCK8 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK8 --check-prefix CK8-64
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK8 --check-prefix CK8-32
+// RUN: %clang_cc1 -DCK8 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK8 --check-prefix CK8-32
+
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK8 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK8 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK8 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}#ifdef CK8
+#ifdef CK8
+struct S1 {
+  int i;
+};
+struct S2 {
+  S1 s;
+  struct S2 *ps;
+};
+
+void test_present_modifier(int arg) {
+  S2 *ps1;
+  S2 *ps2;
+
+  // Make sure the struct picks up present even if another element of the struct
+  // doesn't have present.
+
+  // CK8: private unnamed_addr constant [15 x i64]
+
+  // ps1
+  //
+  // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
+  // MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
+  // MEMBER_OF_1=0x1000000000000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1000000000013
+  // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
+  // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010
+  // PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010
+  // PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013
+  //
+  // CK8-SAME: {{^}} [i64 [[#0x1020]], i64 [[#0x1000000000003]],
+  // CK8-SAME: {{^}} i64 [[#0x1000000000013]], i64 [[#0x1000000001003]],
+  // CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]],
+
+  // arg
+  //
+  // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
+  //
+  // CK8-SAME: {{^}} i64 [[#0x1023]],
+
+  // ps2
+  //
+  // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
+  // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x9000000001003
+  // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x9000000001013
+  // MEMBER_OF_9=0x9000000000000 | FROM=0x2 | TO=0x1 = 0x9000000000003
+  // MEMBER_OF_9=0x9000000000000 | PTR_AND_OBJ=0x10 = 0x9000000000010
+  // PTR_AND_OBJ=0x10 = 0x10
+  // PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13
+  //
+  // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x9000000001003]],
+  // CK8-SAME: {{^}} i64 [[#0x9000000001013]], i64 [[#0x9000000000003]],
+  // CK8-SAME: {{^}} i64 [[#0x9000000000010]], i64 [[#0x10]], i64 [[#0x13]]]
+  #pragma omp target data map(tofrom: ps1->s) \
+                          map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \
+                          map(tofrom: ps2->ps->ps->ps->s)
+  {
+    ++(arg);
+  }
+}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK9 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK9 --check-prefix CK9-64
+// RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-64
+// RUN: %clang_cc1 -DCK9 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-32
+// RUN: %clang_cc1 -DCK9 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-32
+
+// RUN: %clang_cc1 -DCK9 -verify -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK9 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK9 -verify -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// RUN: %clang_cc1 -DCK9 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY2 %s
+// SIMD-ONLY2-NOT: {{__kmpc|__tgt}}
+#ifdef CK9
+void test_present_modifier(int arg) {
+  // PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
+  // CK9: private unnamed_addr constant [1 x i64] [i64 [[#0x1023]]]
+  #pragma omp target data map(present,tofrom: arg)
+  {++arg;}
+}
+#endif
 #endif

clang/test/OpenMP/target_enter_data_codegen.cpp

void foo(int arg) {
   // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   // CK1-NOT: __tgt_target_data_end
   #pragma omp target enter data map(always close, to: lb)
   {++arg;}
 }
 #endif
 ///==========================================================================///
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-64
+// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-64
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-32
+// RUN: %clang_cc1 -DCK1A -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1A --check-prefix CK1A-32
+
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK1A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK1A
+
+// CK1A: [[ST:%.+]] = type { i32, double* }
+template <typename T>
+struct ST {
+  T a;
+  double *b;
+};
+
+ST<int> gb;
+double gc[100];
+
+// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021
+// CK1A: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]]
+
+// PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | TO=0x1 = 0x1425
+// CK1A: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1425]]]
+
+// CK1A-LABEL: _Z3fooi
+void foo(int arg) {
+  int la;
+  float lb[arg];
+
+  // Region 00
+  // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:32|64]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+  // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
+  // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+  // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+  // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+  // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+  // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+  // CK1A-NOT: __tgt_target_data_end
+  #pragma omp target enter data map(present, to: lb)
+  {++arg;}
+
+  // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+  {++arg;}
+
+  // Region 01
+  // CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
+
+  // CK1A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK1A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
+  // CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
+  // CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
+  // CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
+  // CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
+  // CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
+  // CK1A: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
+  // CK1A-NOT: __tgt_target_data_end
+  #pragma omp target enter data map(always close present, to: lb)
+  {++arg;}
+}
+#endif
+///==========================================================================///
 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-64
 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK2 --check-prefix CK2-32
 
 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY1 %s

clang/test/OpenMP/target_map_codegen.cpp

void explicit_maps_single (int ii){
    a++;
   }
 }
 // CK31: define {{.+}}[[CALL00]]
 // CK31: define {{.+}}[[CALL01]]
 
 #endif
 ///==========================================================================///
+// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefixes=CK31A,CK31A-64,CK31A-USE
+// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31A,CK31A-64,CK31A-USE
+// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31A,CK31A-32,CK31A-USE
+// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31A,CK31A-32,CK31A-USE
+
+// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+
+// RUN: %clang_cc1 -DCK31A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefixes=CK31A,CK31A-64,CK31A-NOUSE
+// RUN: %clang_cc1 -DCK31A -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31A,CK31A-64,CK31A-NOUSE
+// RUN: %clang_cc1 -DCK31A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31A,CK31A-32,CK31A-NOUSE
+// RUN: %clang_cc1 -DCK31A -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31A,CK31A-32,CK31A-NOUSE
+
+// RUN: %clang_cc1 -DCK31A -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31A -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31A -verify -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31A -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
+#ifdef CK31A
+
+// CK31A: [[ST:%.+]] = type { i32, i32 }
+struct ST {
+  int i;
+  int j;
+};
+
+// CK31A-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+//
+// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
+// CK31A: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1020]],
+//
+// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
+// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
+// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x1023
+// CK31A-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]], i64 [[#0x1023]],
+//
+// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
+// MEMBER_OF_5=0x5000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x5000000001003
+// MEMBER_OF_5=0x5000000000000 | FROM=0x2 | TO=0x1 = 0x5000000000003
+// CK31A-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x5000000001003]], i64 [[#0x5000000000003]]]
+
+// CK31A-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31A: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]
+//
+// PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | FROM=0x2 | TO=0x1 = 0x1427
+// CK31A: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 [[#0x1427]]]
+
+// CK31A-LABEL: explicit_maps_single{{.*}}(
+void explicit_maps_single (int ii){
+  // CK31A: alloca i32
+
+  // Map of a scalar.
+  // CK31A: [[A:%.+]] = alloca i32
+  int a = ii;
+
+  // CK31A: [[ST1:%.+]] = alloca [[ST]]
+  // CK31A: [[ST2:%.+]] = alloca [[ST]]
+  struct ST st1;
+  struct ST st2;
+
+  // Make sure the struct picks up present even if another element of the struct
+  // doesn't have present.
+  // Region 00
+  // CK31A: [[ST1_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 0
+  // CK31A: [[ST1_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST1]], i{{.+}} 0, i{{.+}} 1
+  // CK31A: [[ST2_I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 0
+  // CK31A: [[ST2_J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[ST2]], i{{.+}} 0, i{{.+}} 1
+  // CK31A-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 7, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[7 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK31A-DAG: [[GEPS]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
+  // CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // st1
+  // CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+  // CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+  // CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP0]]
+  // CK31A-DAG: store i32* [[ST1_I]], i32** [[CP0]]
+  // CK31A-DAG: store i64 %{{.+}}, i64* [[S0]]
+
+  // st1.i
+  // CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK31A-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+  // CK31A-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+  // CK31A-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
+  // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP1]]
+  // CK31A-DAG: store i32* [[ST1_I]], i32** [[CP1]]
+  // CK31A-DAG: store i64 4, i64* [[S1]]
+
+  // st1.j
+  // CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+  // CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+  // CK31A-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+  // CK31A-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+  // CK31A-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
+  // CK31A-DAG: store [[ST]]* [[ST1]], [[ST]]** [[CBP2]]
+  // CK31A-DAG: store i32* [[ST1_J]], i32** [[CP2]]
+  // CK31A-DAG: store i64 4, i64* [[S2]]
+
+  // a
+  // CK31A-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
+  // CK31A-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
+  // CK31A-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
+  // CK31A-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32**
+  // CK31A-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32**
+  // CK31A-DAG: store i32* [[A]], i32** [[CBP3]]
+  // CK31A-DAG: store i32* [[A]], i32** [[CP3]]
+  // CK31A-DAG: store i64 4, i64* [[S3]]
+
+  // st2
+  // CK31A-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4
+  // CK31A-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4
+  // CK31A-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4
+  // CK31A-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to [[ST]]**
+  // CK31A-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i32**
+  // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP4]]
+  // CK31A-DAG: store i32* [[ST2_I]], i32** [[CP4]]
+  // CK31A-DAG: store i64 %{{.+}}, i64* [[S4]]
+
+  // st2.i
+  // CK31A-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5
+  // CK31A-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5
+  // CK31A-DAG: [[S5:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 5
+  // CK31A-DAG: [[CBP5:%.+]] = bitcast i8** [[BP5]] to [[ST]]**
+  // CK31A-DAG: [[CP5:%.+]] = bitcast i8** [[P5]] to i32**
+  // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP5]]
+  // CK31A-DAG: store i32* [[ST2_I]], i32** [[CP5]]
+  // CK31A-DAG: store i64 4, i64* [[S5]]
+
+  // st2.j
+  // CK31A-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6
+  // CK31A-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6
+  // CK31A-DAG: [[S6:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 6
+  // CK31A-DAG: [[CBP6:%.+]] = bitcast i8** [[BP6]] to [[ST]]**
+  // CK31A-DAG: [[CP6:%.+]] = bitcast i8** [[P6]] to i32**
+  // CK31A-DAG: store [[ST]]* [[ST2]], [[ST]]** [[CBP6]]
+  // CK31A-DAG: store i32* [[ST2_J]], i32** [[CP6]]
+  // CK31A-DAG: store i64 4, i64* [[S6]]
+
+  // CK31A-USE: call void [[CALL00:@.+]]([[ST]]* [[ST1]], i32* [[A]], [[ST]]* [[ST2]])
+  // CK31A-NOUSE: call void [[CALL00:@.+]]()
+  #pragma omp target map(tofrom: st1.i) map(present, tofrom: a, st1.j, st2.i) map(tofrom: st2.j)
+  {
+#ifdef USE
+    st1.i++;
+    a++;
+    st1.j++;
+    st2.i++;
+    st2.j++;
+#endif
+  }
+
+  // Always Close Present.
+  // Region 01
+  // CK31A-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
+  // CK31A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK31A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK31A-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK31A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK31A-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]]
+  // CK31A-DAG: store i32* [[VAR0]], i32** [[CP0]]
+
+  // CK31A-USE: call void [[CALL01:@.+]](i32* {{[^,]+}})
+  // CK31A-NOUSE: call void [[CALL01:@.+]]()
+  #pragma omp target map(always close present tofrom: a)
+  {
+#ifdef USE
+    a++;
+#endif
+  }
+}
+// CK31A: define {{.+}}[[CALL00]]
+// CK31A: define {{.+}}[[CALL01]]
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DUSE -DCK31B -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefixes=CK31B,CK31B-64,CK31B-USE
+// RUN: %clang_cc1 -DUSE -DCK31B -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31B,CK31B-64,CK31B-USE
+// RUN: %clang_cc1 -DUSE -DCK31B -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31B,CK31B-32,CK31B-USE
+// RUN: %clang_cc1 -DUSE -DCK31B -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31B,CK31B-32,CK31B-USE
+
+// RUN: %clang_cc1 -DUSE -DCK31B -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31B -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31B -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DUSE -DCK31B -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+
+// RUN: %clang_cc1 -DCK31B -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s --check-prefixes=CK31B,CK31B-64,CK31B-NOUSE
+// RUN: %clang_cc1 -DCK31B -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31B,CK31B-64,CK31B-NOUSE
+// RUN: %clang_cc1 -DCK31B -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31B,CK31B-32,CK31B-NOUSE
+// RUN: %clang_cc1 -DCK31B -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  %s  --check-prefixes=CK31B,CK31B-32,CK31B-NOUSE
+
+// RUN: %clang_cc1 -DCK31B -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31B -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31B -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// RUN: %clang_cc1 -DCK31B -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap  --check-prefix SIMD-ONLY18 %s
+// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
+#ifdef CK31B
+
+// CK31B: [[ST:%.+]] = type { i32, i32 }
+
+// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020
+// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003
+// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003
+
+// CK31B-LABEL: @.__omp_offloading_{{.*}}test_present_members{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
+// CK31B: [[MTYPE00:@.+]] = private {{.*}}constant [3 x i64] [i64 [[#0x1020]],
+// CK31B-SAME: {{^}} i64 [[#0x1000000000003]], i64 [[#0x1000000001003]]]
+
+struct ST {
+  int i;
+  int j;
+  // CK31B-LABEL: define {{.*}}test_present_members{{.*}}(
+  void test_present_members() {
+    // Make sure the struct picks up present even if another element of the
+    // struct doesn't have present.
+    // Region 00
+    // CK31B: [[I:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS:%.+]], i{{.+}} 0, i{{.+}} 0
+    // CK31B: [[J:%.+]] = getelementptr inbounds [[ST]], [[ST]]* [[THIS]], i{{.+}} 0, i{{.+}} 1
+    // CK31B-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 3, i8** [[GEPBP:%[0-9]+]], i8** [[GEPP:%[0-9]+]], i64* [[GEPS:%.+]], i64* getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}})
+    // CK31B-DAG: [[GEPS]] = getelementptr inbounds [3 x i64], [3 x i64]* [[S:%.+]], i{{.+}} 0, i{{.+}} 0
+    // CK31B-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK31B-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+    // this
+    // CK31B-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK31B-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK31B-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
+    // CK31B-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
+    // CK31B-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
+    // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP0]]
+    // CK31B-DAG: store i32* [[I]], i32** [[CP0]]
+    // CK31B-DAG: store i64 %{{.+}}, i64* [[S0]]
+
+    // i
+    // CK31B-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+    // CK31B-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+    // CK31B-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
+    // CK31B-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]**
+    // CK31B-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32**
+    // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP1]]
+    // CK31B-DAG: store i32* [[I]], i32** [[CP1]]
+    // CK31B-DAG: store i64 4, i64* [[S1]]
+
+    // j
+    // CK31B-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+    // CK31B-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+    // CK31B-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
+    // CK31B-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]**
+    // CK31B-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32**
+    // CK31B-DAG: store [[ST]]* [[THIS]], [[ST]]** [[CBP2]]
+    // CK31B-DAG: store i32* [[J]], i32** [[CP2]]
+    // CK31B-DAG: store i64 4, i64* [[S2]]
+
+    // CK31B-USE: call void [[CALL00:@.+]]([[ST]]* [[THIS]])
+    // CK31B-NOUSE: call void [[CALL00:@.+]]()
+    #pragma omp target map(tofrom: i) map(present, tofrom: j)
+    {
+#ifdef USE
+      i++;
+      j++;
+#endif
+    }
+  }
+};
+
+void test() {
+  ST s;
+  s.test_present_members();
+}
+
+// CK31B: define {{.+}}[[CALL00]]
+
+#endif
+///==========================================================================///
 // RUN: %clang_cc1 -DCK32 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK32 --check-prefix CK32-64
 // RUN: %clang_cc1 -DCK32 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK32 --check-prefix CK32-64
 // RUN: %clang_cc1 -DCK32 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK32 --check-prefix CK32-32
 // RUN: %clang_cc1 -DCK32 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK32 --check-prefix CK32-32
 
 // RUN: %clang_cc1 -DCK32 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY32 %s

clang/test/OpenMP/target_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
 // RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp -Wuninitialized
 
-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
 // RUN: %clang_cc1 -DCCODE -verify -fopenmp-simd -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized
 #ifdef CCODE
 void foo(int arg) {
   const int n = 0;
 
   double marr[10][10][10];
 
   #pragma omp target map(marr[2][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
#pragma omp target
       a = 0.0;
       func(arg);
       bf = 20;
     }
     #pragma omp target map(arg,a,d,sref.b)
     {}
     #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) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected addressable lvalue in 'map' clause}}
+    #pragma omp target map(arg,a*2) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}}
     {}
-    #pragma omp target map(arg,(c+1)[2]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    #pragma omp target map(arg,(c+1)[2]) // lt50-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}}
     {}
void func(int arg) {
     #pragma omp target map(close, tofrom: c,f[1:2])
     {}
     #pragma omp target map(close, tofrom: c[:],f)   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
     {}
     #pragma omp target map(close, tofrom: c,f[:])   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
     {}
     #pragma omp target map(close)   // expected-error {{use of undeclared identifier 'close'}}
     {}
+    // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    #pragma omp target map(present, tofrom: c,f)
+    {}
+    // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    #pragma omp target map(present, tofrom: c[1:2],f)
+    {}
+    // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    #pragma omp target map(present, tofrom: c,f[1:2])
+    {}
+    // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    #pragma omp target map(present, tofrom: c[:],f)
+    {}
+    // expected-error@+2 {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    #pragma omp target map(present, tofrom: c,f[:])
+    {}
+    // expected-error@+1 {{use of undeclared identifier 'present'}}
+    #pragma omp target map(present)
+    {}
     #pragma omp target map(close, close, tofrom: a)   // expected-error {{same map type modifier has been specified more than once}}
     {}
-    #pragma omp target map(always, close, always, close, tofrom: a)   // expected-error {{same map type modifier has been specified more than once}} expected-error {{same map type modifier has been specified more than once}}
+    #pragma omp target map(always, close, always, close, tofrom: a)   // expected-error 2 {{same map type modifier has been specified more than once}}
+    {}
+    // ge51-error@+2 {{same map type modifier has been specified more than once}}
+    // lt51-error@+1 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    #pragma omp target map(present, present, tofrom: a)
+    {}
+    // expected-error@+3 2 {{same map type modifier has been specified more than once}}
+    // ge51-error@+2 1 {{same map type modifier has been specified more than once}}
+    // lt51-error@+1 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    #pragma omp target map(always, close, present, always, close, present, tofrom: a)
     {}
     #pragma omp target map( , tofrom: a)   // expected-error {{missing map type modifier}}
     {}
     #pragma omp target map( , , tofrom: a)   // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}}
     {}
     #pragma omp target map( , , : a)   // expected-error {{missing map type modifier}} expected-error {{missing map type modifier}} expected-error {{missing map type}}
     {}
-    #pragma omp target map( d, f, bf: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
-    {}
-    #pragma omp target map( , f, : a)   // expected-error {{missing map type modifier}} expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+    // ge51-error@+3 2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+    // lt51-error@+2 2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    // expected-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+    #pragma omp target map( d, f, bf: a)
+    {}
+    // expected-error@+4 {{missing map type modifier}}
+    // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+    // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    // expected-error@+1 {{missing map type}}
+    #pragma omp target map( , f, : a)
     {}
     #pragma omp target map(always close: a)   // expected-error {{missing map type}}
     {}
     #pragma omp target map(always close bf: a)   // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
     {}
-    #pragma omp target map(always tofrom close: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
-    {}
-    #pragma omp target map(tofrom from: a)   // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    // ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+    // lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    // expected-error@+1 {{missing map type}}
+    #pragma omp target map(always tofrom close: a)
+    {}
+    // ge51-error@+2 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+    // lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+    #pragma omp target map(tofrom from: a)
     {}
     #pragma omp target map(close bf: a)   // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
     {}
-    #pragma omp target map(([b[I]][bf])f)  // le45-error {{expected ',' or ']' in lambda capture list}} le45-error {{expected ')'}} le45-note {{to match this '('}}
+    #pragma omp target map(([b[I]][bf])f)  // lt50-error {{expected ',' or ']' in lambda capture list}} lt50-error {{expected ')'}} lt50-note {{to match this '('}}
     {}
     return;
   }
 };
 
 struct SB {
   unsigned A;
   unsigned B;

 #pragma omp target map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
   {}
 #pragma omp target map(r.RPtrS, r.RPtrS->B)    // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
   {}
 #pragma omp target map(r.RPtrS->A, r.RPtrS->B)
   {}
 #pragma omp target map(r.S.Arr[:12])
   {}
-#pragma omp target map(r.S.foo() [:12]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected addressable lvalue in 'map' clause}}
+#pragma omp target map(r.S.foo() [:12]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}}
   {}
 #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.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) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map((p + 1)->A) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
 #pragma omp target map(u.B)        // expected-error {{mapping of union members is not allowed}}
   {}
 #pragma omp target
   {
     u.B = 0;
     r.S.foo();
   }
T tmain(T argc) {
   const T da[5] = { 0 };
   S4 e(4);
   S5 g(5);
   T i, t[20];
   T &j = i;
   T *k = &j;
   T x;
   T y;
-  T to, tofrom, always, close;
+  T to, tofrom, always, close, present;
   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(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'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   foo();
 #pragma omp target map(T) // expected-error {{'T' does not refer to a value}}
   foo();
-#pragma omp target map(I) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error 2 {{expected addressable lvalue in 'map' clause}}
+#pragma omp target map(I) // lt50-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error 2 {{expected addressable lvalue in 'map' clause}}
   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 data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
-#pragma omp target data map(tofrom: argc > 0 ? x : y) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error 2 {{expected addressable lvalue in 'map' clause}}
+#pragma omp target data map(tofrom: argc > 0 ? x : y) // lt50-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error 2 {{expected addressable lvalue in 'map' clause}}
 #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}} warn-warning 2 {{Type 'const S2' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning 2 {{Type 'const S3' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(ba) // warn-warning 2 {{Type 'const S2 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(ca) // warn-warning 2 {{Type 'const S3 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #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(k[:4], j, l[:5]) // expected-note 2 {{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)
   foo();
 
 #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', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target data map(tofrom, always: x)
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
 #pragma omp target data map(close, tofrom: x)
 #pragma omp target data map(close: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target data map(tofrom, close: x)
 #pragma omp target data map(close, tofrom: close, tofrom, x)
   foo();
 
+// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+#pragma omp target data map(present, tofrom: x)
+// ge51-error@+2 {{missing map type}}
+// lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target data map(present: x)
+// ge51-error@+4 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+3 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// ge51-error@+2 {{missing map type}}
+// lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target data map(tofrom, present: x)
+// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+#pragma omp target data map(present, tofrom: present, tofrom, x)
+  foo();
+
   T marr[10][10], iarr[5];
 #pragma omp target data map(marr[10][0:2:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
   {}
 #pragma omp target data map(iarr[:2:d]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
   {}
   return 0;
 }
 
int main(int argc, char **argv) {
   S4 e(4);
   S5 g(5);
   int i;
   int &j = i;
   int *k = &j;
   S6<int> m;
   int x;
   int y;
-  int to, tofrom, always, close;
+  int to, tofrom, always, close, present;
   const int (&l)[5] = da;
   SC1 s;
   SC1 *p;
   int Arr[10];
-#pragma omp target data map // expected-error {{expected '(' after 'map'}} le45-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} le50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
+#pragma omp target data map // expected-error {{expected '(' after 'map'}} lt50-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} ge50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
 #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 data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
-#pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected addressable lvalue in 'map' clause}}
+#pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2]) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}} ge50-error {{expected addressable lvalue in 'map' clause}}
 #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}} warn-warning {{Type 'const S2' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'const S3' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(argv[1])
 #pragma omp target data map(ba) // warn-warning {{Type 'const S2 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(ca) // warn-warning {{Type 'const S3 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note {{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)
   foo();
 
 #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', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target data map(tofrom, always: x)
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 #pragma omp target data map(close, tofrom: x)
 #pragma omp target data map(close: x) // expected-error {{missing map type}}
-#pragma omp target data map(tofrom, close: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target data map(tofrom, close: x)
+  foo();
+// lt51-error@+1 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+#pragma omp target data map(present, tofrom: x)
+// ge51-error@+2 {{missing map type}}
+// lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target data map(present: x)
+// ge51-error@+4 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+3 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// ge51-error@+2 {{missing map type}}
+// lt51-error@+1 {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target data map(tofrom, present: x)
   foo();
 #pragma omp target private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target' directive}}  expected-note {{defined as private}}
   {}
 #pragma omp target firstprivate(j) map(j)  // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target' directive}} expected-note {{defined as firstprivate}}
   {}
 #pragma omp target map(m) // warn-warning {{Type 'S6<int>' is not trivially copyable and not guaranteed to be mapped correctly}}
   {}
 // expected-note@+1 {{used here}}
// expected-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
 #pragma omp target map(s.s.s.b[:2])
   { s.s.s.b[0]++; }
 #pragma omp target map(s8[0:1], s9) // warn-warning {{Type 'class S8' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'class S9' is not trivially copyable and not guaranteed to be mapped correctly}}
   {}
 
   int **BB, *offset, *a;
 
-#pragma omp target map(**(BB+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(BB+*offset)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
-#pragma omp target map(**(BB+y)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(BB+y)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
-#pragma omp target map(*(a+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(*(a+*offset)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
-#pragma omp target map(**(*offset+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(*offset+BB)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
-#pragma omp target map(**(y+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(y+BB)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
-#pragma omp target map(*(*offset+a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(*(*offset+a)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
-#pragma omp target map(**(*offset+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(*offset+BB+*a)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
-#pragma omp target map(**(*(*(&offset))+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target map(**(*(*(&offset))+BB+*a)) // lt50-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
 #pragma omp target map(*(a+(a))) // expected-error {{invalid operands to binary expression ('int *' and 'int *')}}
   {}
 #pragma omp target map(*(1+*a+*a)) // expected-error {{indirection requires pointer operand ('int' invalid)}}
   {}
 
 #pragma omp target map(delete: a) // expected-error {{map type 'delete' is not allowed for '#pragma omp target'}}
   {}
 #pragma omp target map(release: a) // expected-error {{map type 'release' is not allowed for '#pragma omp target'}}
   {}
 
   int marr[10][10], iarr[5];
 
 #pragma omp target map(marr[10][0:2:2]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
   {}
 #pragma omp target map(iarr[:2:d]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
   {}
 
-#pragma omp target data map(Arr[0:4]) // le45-note {{used here}}
+#pragma omp target data map(Arr[0:4]) // lt50-note {{used here}}
   {
 #pragma omp target
-    Arr[0] = 2; // le45-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
+    Arr[0] = 2; // lt50-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
   }
 
   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}}
 }
 #endif

clang/test/OpenMP/target_parallel_for_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-
-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 #pragma omp target parallel for map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(tofrom: t[:I])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(T) // expected-error {{'T' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(I) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target parallel for map(I)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
+// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target parallel for map(tofrom \
-                                    : argc > 0 ? x : y) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+                                    : argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(ba)

 #pragma omp target parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target parallel for map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }
#pragma omp target parallel for map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i)
     foo();
+// ge50-error@+3 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target parallel for map(tofrom \
-                                    : argc > 0 ? argv[1] : argv[2]) // omp4-error {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error {{expected addressable lvalue in 'map' clause}}
+                                    : argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(argv[1])

 #pragma omp target parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target parallel for map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target parallel for'}}
   for (i = 0; i < argc; ++i)
     foo();
 #pragma omp target parallel for map(release: j) // expected-error {{map type 'release' is not allowed for '#pragma omp target parallel for'}}
   for (i = 0; i < argc; ++i)
     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}}
 }
 

clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
-
-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp-simd -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp-simd -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 %s -Wno-openmp-mapping -Wuninitialized
+
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 #pragma omp target parallel for simd map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(tofrom: t[:I])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(T) // expected-error {{'T' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(I) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target parallel for simd map(I)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(x, z)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
+// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target parallel for simd map(tofrom \
-                                         : argc > 0 ? x : y) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+                                         : argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(ba)

 #pragma omp target parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target parallel for simd map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }
#pragma omp target parallel for simd map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to, z)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i)
     foo();
+// ge50-error@+3 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target parallel for simd map(tofrom \
-                                         : argc > 0 ? argv[1] : argv[2]) // omp4-error {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error {{expected addressable lvalue in 'map' clause}}
+                                         : argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(argv[1])

 #pragma omp target parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target parallel for simd map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target parallel for simd'}}
   for (i = 0; i < argc; ++i)
     foo();
 #pragma omp target parallel for simd map(release: j) // expected-error {{map type 'release' is not allowed for '#pragma omp target parallel for simd'}}
   for (i = 0; i < argc; ++i)
     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}}
 }
 

clang/test/OpenMP/target_parallel_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-
-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 #pragma omp target parallel map(x, z)
   foo();
 #pragma omp target parallel map(tofrom: t[:I])
   foo();
 #pragma omp target parallel map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   foo();
 #pragma omp target parallel map(T) // expected-error {{'T' does not refer to a value}}
   foo();
-#pragma omp target parallel map(I) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target parallel map(I)
   foo();
 #pragma omp target parallel map(S2::S2s)
   foo();
 #pragma omp target parallel map(S2::S2sc)
   foo();
 #pragma omp target parallel map(x)
   foo();
 #pragma omp target parallel map(to: x)
   foo();
 #pragma omp target parallel map(to: to)
   foo();
 #pragma omp target parallel map(to)
   foo();
 #pragma omp target parallel map(to, x)
   foo();
 #pragma omp target parallel map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
+// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target parallel map(tofrom \
-                                : argc > 0 ? x : y) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+                                : argc > 0 ? x : y)
   foo();
 #pragma omp target parallel map(argc)
   foo();
 #pragma omp target parallel map(S1) // expected-error {{'S1' does not refer to a value}}
   foo();
 #pragma omp target parallel map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   foo();
 #pragma omp target parallel map(ba)

 #pragma omp target parallel map(l)
   foo();
 }
 
 #pragma omp target parallel map(always, tofrom: x)
   foo();
 #pragma omp target parallel map(always: x) // expected-error {{missing map type}}
   foo();
-#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target parallel map(tofrom, always: x)
   foo();
 #pragma omp target parallel map(always, tofrom: always, tofrom, x)
   foo();
 #pragma omp target parallel map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
   return 0;
 }

 #pragma omp target parallel map(to: to)
   foo();
 #pragma omp target parallel map(to)
   foo();
 #pragma omp target parallel map(to, x)
   foo();
 #pragma omp target parallel map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
+// ge50-error@+3 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target parallel map(tofrom \
-                                : argc > 0 ? argv[1] : argv[2]) // omp4-error {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error {{expected addressable lvalue in 'map' clause}}
+                                : argc > 0 ? argv[1] : argv[2])
   foo();
 #pragma omp target parallel map(argc)
   foo();
 #pragma omp target parallel map(S1) // expected-error {{'S1' does not refer to a value}}
   foo();
 #pragma omp target parallel map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   foo();
 #pragma omp target parallel map(argv[1])

 #pragma omp target parallel map(l)
   foo();
 }
 
 #pragma omp target parallel map(always, tofrom: x)
   foo();
 #pragma omp target parallel map(always: x) // expected-error {{missing map type}}
   foo();
-#pragma omp target parallel map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target parallel map(tofrom, always: x)
   foo();
 #pragma omp target parallel map(always, tofrom: always, tofrom, x)
   foo();
 #pragma omp target parallel map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 #pragma omp target parallel map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target parallel'}}
   foo();
 #pragma omp target parallel map(release: j) // expected-error {{map type 'release' is not allowed for '#pragma omp target parallel'}}
   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}}
 }
 

clang/test/OpenMP/target_simd_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
-
-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp-simd -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp-simd -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 %s -Wno-openmp-mapping -Wuninitialized
+
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 #pragma omp target simd map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(tofrom: t[:I])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(T) // expected-error {{'T' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(I) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target simd map(I)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
+// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target simd map(tofrom \
-                            : argc > 0 ? x : y) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+                            : argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(ba)

 #pragma omp target simd map(l) // OK
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target simd map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }
#pragma omp target simd map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i)
     foo();
+// ge50-error@+3 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target simd map(tofrom \
-                            : argc > 0 ? argv[1] : argv[2]) // omp4-error {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error {{expected addressable lvalue in 'map' clause}}
+                            : argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(argv[1])

 #pragma omp target simd map(l) //
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target simd map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target simd'}}
   for (i = 0; i < argc; ++i)
     foo();
 #pragma omp target simd map(release: j) // expected-error {{map type 'release' is not allowed for '#pragma omp target simd'}}
   for (i = 0; i < argc; ++i)
     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}}
 }
 

clang/test/OpenMP/target_teams_distribute_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-
-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 #pragma omp target teams distribute map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(tofrom: t[:I])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(T) // expected-error {{'T' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(I) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams distribute map(I)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
+// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target teams distribute map(tofrom \
-                                        : argc > 0 ? x : y) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+                                        : argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(ba)

 #pragma omp target teams distribute map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target teams distribute map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }
#pragma omp target teams distribute map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i)
     foo();
+// ge50-error@+3 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target teams distribute map(tofrom \
-                                        : argc > 0 ? argv[1] : argv[2]) // omp4-error {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error {{expected addressable lvalue in 'map' clause}}
+                                        : argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(argv[1])

 #pragma omp target teams distribute map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target teams distribute map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target teams distribute'}}
   for (i = 0; i < argc; ++i)
     foo();
 #pragma omp target teams distribute map(release: j) // expected-error {{map type 'release' is not allowed for '#pragma omp target teams distribute'}}
   for (i = 0; i < argc; ++i)
     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}}
 }
 

clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 #pragma omp target teams distribute parallel for map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(tofrom: t[:I])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(T) // expected-error {{'T' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(I) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams distribute parallel for map(I)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(x, z)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
+// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target teams distribute parallel for map(tofrom \
-                                                     : argc > 0 ? x : y) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+                                                     : argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(ba)

 #pragma omp target teams distribute parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target teams distribute parallel for map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute parallel for map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
+// ge50-error@+3 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target teams distribute parallel for map(tofrom \
-                                                     : argc > 0 ? argv[1] : argv[2]) // omp4-error {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error {{expected addressable lvalue in 'map' clause}}
+                                                     : argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(argv[1])

 #pragma omp target teams distribute parallel for map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target teams distribute parallel for map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target teams distribute parallel for'}}
   for (i = 0; i < argc; ++i)
     foo();
 #pragma omp target teams distribute parallel for map(release: j) // expected-error {{map type 'release' is not allowed for '#pragma omp target teams distribute parallel for'}}
   for (i = 0; i < argc; ++i)
     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}}
 }
 

clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-
-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 #pragma omp target teams distribute parallel for simd map(x, z)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(tofrom: t[:I])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(T) // expected-error {{'T' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(I) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams distribute parallel for simd map(I)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
+// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target teams distribute parallel for simd map(tofrom \
-                                                          : argc > 0 ? x : y) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+                                                          : argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(ba)

 #pragma omp target teams distribute parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target teams distribute parallel for simd map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

 #pragma omp target teams distribute parallel for simd map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to, x, z)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
+// ge50-error@+3 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target teams distribute parallel for simd map(tofrom \
-                                                          : argc > 0 ? argv[1] : argv[2]) // omp4-error {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error {{expected addressable lvalue in 'map' clause}}
+                                                          : argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(argv[1])

 #pragma omp target teams distribute parallel for simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute parallel for simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target teams distribute parallel for simd map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target teams distribute parallel for simd'}}
   for (i = 0; i < argc; ++i)
     foo();
 #pragma omp target teams distribute parallel for simd map(release: j) // expected-error {{map type 'release' is not allowed for '#pragma omp target teams distribute parallel for simd'}}
   for (i = 0; i < argc; ++i)
     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}}
 }
 

clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-
-// RUN: %clang_cc1 -verify=expected,omp4 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+
+// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 #pragma omp target teams distribute simd map(x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(tofrom: t[:I])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(T) // expected-error {{'T' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(I) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams distribute simd map(I)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(x, z)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
+// ge50-error@+3 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target teams distribute simd map(tofrom \
-                                             : argc > 0 ? x : y) // omp4-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error 2 {{expected addressable lvalue in 'map' clause}}
+                                             : argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(ba)

 #pragma omp target teams distribute simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target teams distribute simd map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }
#pragma omp target teams distribute simd map(to: to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i)
     foo();
+// ge50-error@+3 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target teams distribute simd map(tofrom \
-                                             : argc > 0 ? argv[1] : argv[2]) // omp4-error {{expected expression containing only member accesses and/or array sections based on named variables}} omp5-error {{expected addressable lvalue in 'map' clause}}
+                                             : argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(argc)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(S1) // expected-error {{'S1' does not refer to a value}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(argv[1])

 #pragma omp target teams distribute simd map(l)
   for (i = 0; i < argc; ++i) foo();
 }
 
 #pragma omp target teams distribute simd map(always, tofrom: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always: x) // expected-error {{missing map type}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target teams distribute simd map(tofrom, always: x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(always, tofrom: always, tofrom, x)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) 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}}
 }
 

clang/test/OpenMP/target_teams_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-version=50 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,lt45,lt50,lt51 -fopenmp-version=40 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge45,lt50,lt51 -fopenmp-version=45 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp-version=50 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge45,ge50,ge51 -fopenmp-version=51 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
 
-// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
-// RUN: %clang_cc1 -DCCODE -verify=expected,le45 -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,ge45,ge50,lt51 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -DCCODE -verify=expected,ge45,ge50,lt51 -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-mapping -Wuninitialized
 #ifdef CCODE
 void foo(int arg) {
   const int n = 0;
 
   double marr[10][10][10];
 
   #pragma omp target teams map(marr[2][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
   {}
struct SA {
   T d;
   float e[I];
   T *f;
   void func(int arg) {
     #pragma omp target teams map(arg,a,d)
     {}
     #pragma omp target teams map(arg[2:2],a,d) // expected-error {{subscripted value is not an array or pointer}}
     {}
-    #pragma omp target teams map(arg,a*2) // le50-error {{expected addressable lvalue in 'map' clause}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    // ge50-error@+2 {{expected addressable lvalue in 'map' clause}}
+    // lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+    #pragma omp target teams map(arg,a*2)
     {}
-    #pragma omp target teams map(arg,(c+1)[2]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    // lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+    #pragma omp target teams map(arg,(c+1)[2])
     {}
     #pragma omp target teams map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}}
     {}
     #pragma omp target teams map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}}
     {}
 
     #pragma omp target teams map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
     {}
void SAclient(int arg) {
   #pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer dereferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
   {}
   #pragma omp target teams map(r.RPtrS, r.RPtrS->B) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
   {}
   #pragma omp target teams map(r.RPtrS->A, r.RPtrS->B)
   {}
   #pragma omp target teams map(r.S.Arr[:12])
   {}
-  #pragma omp target teams map(r.S.foo()[:12]) // le50-error {{expected addressable lvalue in 'map' clause}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  // ge50-error@+2 {{expected addressable lvalue in 'map' clause}}
+  // lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+  #pragma omp target teams map(r.S.foo()[:12])
   {}
   #pragma omp target teams map(r.C, r.D)
   {}
   #pragma omp target teams map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
   {}
   #pragma omp target teams map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
   {}
   #pragma omp target teams 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.
void SAclient(int arg) {
   #pragma omp target teams map(r.ArrS[3].Arr[6:7])
   {}
   #pragma omp target teams map(r.S.Arr[4:5])
   {}
   #pragma omp target teams map(r.S.Ptr[4:5])
   {}
   #pragma omp target teams map(r.S.Ptr[:])  // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
   {}
-  #pragma omp target teams map((p+1)->A) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  // lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+  #pragma omp target teams map((p+1)->A)
   {}
   #pragma omp target teams map(u.B)  // expected-error {{mapping of union members is not allowed}}
   {}
 
   #pragma omp target data map(to: r.C) //expected-note {{used here}}
   {
     #pragma omp target teams 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 teams map(x)
   foo();
 #pragma omp target teams map(tofrom: t[:I])
   foo();
 #pragma omp target teams map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
   foo();
 #pragma omp target teams map(T) // expected-error {{'T' does not refer to a value}}
   foo();
-#pragma omp target teams map(I) // le50-error 2 {{expected addressable lvalue in 'map' clause}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(I)
   foo();
 #pragma omp target teams map(S2::S2s)
   foo();
 #pragma omp target teams map(S2::S2sc)
   foo();
 #pragma omp target teams map(x)
   foo();
 #pragma omp target teams map(to: x)
   foo();
 #pragma omp target teams map(to: to)
   foo();
 #pragma omp target teams map(to)
   foo();
 #pragma omp target teams map(to, x)
   foo();
 
 #pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
-#pragma omp target data map(tofrom: argc > 0 ? x : y) // le50-error 2 {{expected addressable lvalue in 'map' clause}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+// ge50-error@+2 2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target data map(tofrom: argc > 0 ? x : y)
 
 #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}}
 #pragma omp target data map(ba)
 #pragma omp target data map(ca)
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{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 teams map(l)
   foo();
 
 #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', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target data map(tofrom, always: x)
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target teams 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;
   S6<int> m;
   int x;
   int y;
   int to, tofrom, always;
   const int (&l)[5] = da;
-#pragma omp target data map // expected-error {{expected '(' after 'map'}} le45-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} le50-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
+// expected-error@+3 {{expected '(' after 'map'}}
+// ge50-error@+2 {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
+// lt50-error@+1 {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}}
+#pragma omp target data 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 teams map(x)
   foo();
 
 #pragma omp target teams map(to: x)
   foo();
 #pragma omp target teams map(to: to)
   foo();
 #pragma omp target teams map(to)
   foo();
 #pragma omp target teams map(to, x)
   foo();
 
 #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]) // le50-error {{expected addressable lvalue in 'map' clause}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+// ge50-error@+2 {{expected addressable lvalue in 'map' clause}}
+// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2])
 #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}}
 #pragma omp target data map(argv[1])
 #pragma omp target data map(ba)
 #pragma omp target data map(ca)
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)

 #pragma omp target data map(k[:4], j, l[:5]) // expected-note {{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 teams map(l)
   foo();
 
 #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', 'close', or 'mapper'}} expected-error {{missing map type}}
+// ge51-error@+3 {{incorrect map type modifier, expected 'always', 'close', 'mapper', or 'present'}}
+// lt51-error@+2 {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}}
+// expected-error@+1 {{missing map type}}
+#pragma omp target data map(tofrom, always: x)
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
-#pragma omp target teams private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}}  le45-note {{defined as private}}
+// lt50-error@+2 {{private variable cannot be in a map clause in '#pragma omp target teams' directive}}
+// lt50-note@+1 {{defined as private}}
+#pragma omp target teams private(j) map(j)
   {}
 
-#pragma omp target teams firstprivate(j) map(j) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} le45-note {{defined as firstprivate}}
+// lt50-error@+2 {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}}
+// lt50-note@+1 {{defined as firstprivate}}
+#pragma omp target teams firstprivate(j) map(j)
   {}
 
 #pragma omp target teams map(m)
   {}
 
   int **BB, *offset, *a;
 
-#pragma omp target teams map(**(BB+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(**(BB+*offset))
   {}
-#pragma omp target teams map(**(BB+y)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(**(BB+y))
   {}
-#pragma omp target teams map(*(a+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(*(a+*offset))
   {}
-#pragma omp target teams map(**(*offset+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(**(*offset+BB))
   {}
-#pragma omp target teams map(**(y+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(**(y+BB))
   {}
-#pragma omp target teams map(*(*offset+a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(*(*offset+a))
   {}
-#pragma omp target teams map(**(*offset+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(**(*offset+BB+*a))
   {}
-#pragma omp target teams map(**(*(*(&offset))+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+// lt50-error@+1 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(**(*(*(&offset))+BB+*a))
   {}
 #pragma omp target teams map(*(a+(a))) // expected-error {{invalid operands to binary expression ('int *' and 'int *')}}
   {}
 #pragma omp target teams map(*(1+*a+*a)) // expected-error {{indirection requires pointer operand ('int' invalid)}}
   {}
 
 #pragma omp target teams map(delete: j) // expected-error {{map type 'delete' is not allowed for '#pragma omp target teams'}}
   {}
 #pragma omp target teams map(release: j) // expected-error {{map type 'release' is not allowed for '#pragma omp target teams'}}
   {}
 
   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}}
 }
 #endif