[WIP][OPENMP5.0] allow lvalue for motion clause

[cchen]

Summary

In OpenMP 5.0 we allow any lvalue for motion clause and map clause, below are some examples for lvalue:

int *p;

int **pp;

struct S {
  int struct *sp;
  double *d;
  struct S **sp_arr;
};
struct S *s;

#pragma omp target update to(*p)
#pragma omp target update to(*(p+3))
#pragma omp target update to((p+3)[k])

#pragma omp target update to(**pp)
#pragma omp target update to(*pp[0])
#pragma omp target update to(*(pp+l)[2])
#pragma omp target update to((pp+l)[2][3])

#pragma omp target update to(*(s->d))
#pragma omp target update to(*(s->sp->d))
#pragma omp target update to(*(s->sp->d+k))
#pragma omp target update to((s->sp->d+k)[2])
#pragma omp target update to(*(s->sp->sp_arr+k))

Now only add test for motion clause, after the implementation is fine,
will add test for map since the core logic is in the same helper
function.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7707–7709

This assertion must check that the expression is just an lvalue, no?

clang/lib/Sema/SemaOpenMP.cpp

15525–15526

Why just the LHS is analyzed? Also, what about support for other expressions, like casting, call, choose etc., which may result in lvalue?

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7707–7709

Yes.

clang/lib/Sema/SemaOpenMP.cpp

15525–15526

1. LHS: I'll fix that
2. I'll add support for casting, call, etc
3. For "choose" are you referring to something like (a < b ? b : a)?

[ABataev]

Not sure that we need to dig that deep. Maybe just keep thw original analysis for the known expressions and map all other lvalues without some extra analysis, just like DeclRefExpr nodes?

[ABataev]

Hmm, we still need to find some basic decl to remap it successfully at the codegen. Not sure that we'll be able to support it in full. It would be good to investigate how we can handle them at the codegen.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15525–15526

For the handling of BinaryOperator, I'm not sure why should we handle RHS since the possible source code I can imagine is *(ptr+l) or something like (ptr+l)[3].

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15525–15526

*(2+ptr) is correct too. And, btw, 3[ptr+1] too, especially in C. Moreover, something like *(b+c) is also allowed. That's why I said that better to avoid deep analysis of lvalues, there are too many combinations and better to switch to something basic.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15525–15526

But at least we need the base declaration for motion/map clause, right? And to achieve this, we need to peel the expression to get the DeclRefExpr.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15525–15526

What's the base declaration in *(a+b)?

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15525–15526

I understand your point, I'm just not sure how to handle (do not analyze) this given that we need to fill VarDecl for OMPToClause/OMPFromClause.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15525–15526

But is *(a+b) even valid? Clang and GCC both emit error message: https://godbolt.org/z/jqDBx_

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15525–15526

Yes, in this form it is not. But here is another form *(a+*b) or *(a+(unsigned long long)b). You an try to filter rvalues out of all these forms, though and still try to find a base decl. But still, there are going to be some cases where you won't be possible to handle them correctly, like foo().
If you still want to dig that deep, I would recommend to use RecursiseStmtVisitor. If an expression is rvalue - throw it out of analysis. If base decl is not single, not variable and not member decl - emit an error.

[cchen]

Add a checker for locator (test is not complete yet).

The logic of finding basic decl in LocatorChecker:

• int *B; omp target map(*B)
  • B
    • pointer count (* in type): 1
    • deref count: 1
    • B is basic decl since pointer count equal to deref count
• int *B, l; omp target map(*(B+l))
  • B
    • pointer count (* in type): 1
    • deref count: 1
    • B is basic decl since pointer count equal to deref count
  • l
    • pointer count (* in type): 0
    • deref count: 1
    • pointer count = 0, deref count 1 => not equal
• int **B, **l; omp target map((B+**l)[1][2]
  • B
    • pointer count (* in type): 2
    • deref count: 2 (2 array subscript)
    • B is basic decl since pointer count equal to deref count
  • l
    • pointer count (* in type): 2
    • deref count: 4 (2 array subscript, 2 unary op)
    • pointer count = 2, deref count 4 => not equal

[jdoerfert]

Thanks for working on this! While you are at it, *this is probably one of the most important ones to test and support.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7708–7709

This must be updated, I assume. Instead need to check that it is an lvalue.

7837

Here we should not have any functiondecls. Must be an assert.

clang/lib/Sema/SemaOpenMP.cpp

15212

I don't think anything except for VarDecls must be allowed here.

15235

Need a check that this is a dereference op. Also, maybe allow using an addr_of operation?

15244–15245

Why you have all these comments here?

15267–15271

Some extra checks for the allowed operators are required.

15268

FunctionDecls should not be allowed here

15276

I would add a check that the child is and expression and not a glvalue, then do not analyze it.

15281

getFoundBase or something like this?

15283

Why do you need components list here?

15327

Same here, no function decls

15521

Why are using IgnoreParenImpCasts() here? It drops implicit casts to RValue.

16058–16059

Message should be updated too.

16059

Why MLV_LValueCast must be allowed here?

16124

Again, functiondecls must not be allowed here. I would suggest emitting a warning for them (that user mapped a function) a drop them from the mapping list.

clang/test/OpenMP/target_teams_map_messages.cpp

441–442

Not sure that this must be supported. What is the base decl here, x or y?

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7837

Don't we need to support something like #pragma omp target map(r.S.foo()[:12])?

clang/lib/Sema/SemaOpenMP.cpp

15283

Since I need to insert the information in VisitDeclRefExpr.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7837

It does not lead to the actual mapping. Better to inform the user about this rather than consuming it silently.

clang/lib/Sema/SemaOpenMP.cpp

15283

You can do this outside of this class if you need to do it. Still. you have this GetDRE() member function.

[cchen]

In D72811#1837392, @jdoerfert wrote:

Thanks for working on this! While you are at it, *this is probably one of the most important ones to test and support.

Can anyone tell me how to get ValueDecl from CXXThisExpr? I cannot find any method return decl in https://clang.llvm.org/doxygen/classclang_1_1CXXThisExpr.html#details.
Thanks!

clang/lib/Sema/SemaOpenMP.cpp

15235

is addr_of operation allowed in lvalue?

In this code:

int arr[50];

#pragma omp target map(&arr)
  {}

We now reject &arr since RE->IgnoreParenImpCasts()->isLValue() return false. (RE is the expr of &arr)

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15235

BTW, RE->isLValue() also return false in this case.

[ABataev]

In D72811#1839538, @cchen wrote:

In D72811#1837392, @jdoerfert wrote:

Thanks for working on this! While you are at it, *this is probably one of the most important ones to test and support.

Can anyone tell me how to get ValueDecl from CXXThisExpr? I cannot find any method return decl in https://clang.llvm.org/doxygen/classclang_1_1CXXThisExpr.html#details.
Thanks!

CXXThisExpr is an expression without associated declaration.

clang/lib/Sema/SemaOpenMP.cpp

15235

I'm not saying that &var must be allowed, but something like *(&var). Of course, just &var is an rvalue.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15235

Got it, thanks for the clarification!

[cchen]

In D72811#1839561, @ABataev wrote:

In D72811#1839538, @cchen wrote:

In D72811#1837392, @jdoerfert wrote:

Thanks for working on this! While you are at it, *this is probably one of the most important ones to test and support.

Can anyone tell me how to get ValueDecl from CXXThisExpr? I cannot find any method return decl in https://clang.llvm.org/doxygen/classclang_1_1CXXThisExpr.html#details.
Thanks!

CXXThisExpr is an expression without associated declaration.

Then which kind of thing should I sent to CurComponents if we do not have associated declaration for CXXThisExpr? We can't just send nullptr, right?

clang/lib/Sema/SemaOpenMP.cpp

16059

I'm doing this since isLValue return false for either C style casting CXX style casting and MLV_LValueCast can categorize these two as MLV_LValueCast.

[ABataev]

In D72811#1843066, @cchen wrote:

In D72811#1839561, @ABataev wrote:

In D72811#1839538, @cchen wrote:

In D72811#1837392, @jdoerfert wrote:

Thanks for working on this! While you are at it, *this is probably one of the most important ones to test and support.

Can anyone tell me how to get ValueDecl from CXXThisExpr? I cannot find any method return decl in https://clang.llvm.org/doxygen/classclang_1_1CXXThisExpr.html#details.
Thanks!

CXXThisExpr is an expression without associated declaration.

Then which kind of thing should I sent to CurComponents if we do not have associated declaration for CXXThisExpr? We can't just send nullptr, right?

Check how compiler handles mapping of this[0:1], which is the same as *this.

clang/lib/Sema/SemaOpenMP.cpp

16059

Are they the lvalues according to the standard? If not, no need for this analysis.

[cchen]

Fix based on feedback

[cchen]

Update due to negligence

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15201

Seems to me, you're skipping many expressions, which could be supported. Did you try to test it with casting expressions and others?

15203–15205

Use default initializers for fields.

15213

Better to use early exit where possible to reduce structural complexity of the functions.

15215–15216

This really looks ugly. Need to find a better solution.

15233

Why do you need to count the number of dereferences? And have this comparison here?

15267–15268

Restore original code

15267–15268

If you have the visitor class, you can replace all this code with an analysis in the visitor.

15523

General question about several cases. How we're going to support them?

1. (a ? b : c).
2. __builtin_choose_expr(a, b, c).
3. a = b.
4. a?:b
5. __builtin_convertvector(x, ty)
6. (int&)a
7. v.xy, where v is an extended vector
8. a.b, where b is a bitfield
9. __builtin_bit_cast(v, ty)
10. const_cast<ty &>(a)
11. dynamic_cast<ty &>(a)
12. reinterpret_cast
13. static_cast
14. typeid() (also is an lvalue).
15. __uuidof(*comPtr)
16. lambda calls
17. User defined literals

15524

Better to pass CurComponents as an argument at the construction of the checker to avoid extra filling/copying.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15523

I think we could first evaluate the lvalue, and then pass the result of &lvalue to device?

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15523

What do you mean when you say evaluate lvalue? In veneral, it can be evaluated at the runtime. Do you propose to implement dyic translation? It will definetely slow down the execution on the device.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15523

I mean evaluate lvalue before sending &lvalue to device. For example:

int a, b;
b = 5;
#pragma omp target map(a = b) // assign b to a before sending `&(a=b)` to device
{
   a++;
}

Therefore, the above example have the same semantics as this:

int a, b;
b = 5;
int &c = (a=b)
#pragma omp target map(c)
{
   a++;
}

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15523

Sure, we do this already, generally speaking. The main problem here is how to map the address and associate it with address on the device. We do it at the compile time, support of general lvalues requires runtime translation + special instrumentation of the device code for address translation.

[cchen]

1. Refactor based on feedback and use the visitor class to integrate the analysis

in checkMapClauseExpressionBase.

2. Handle more lvalue cases

[ABataev]

I would suggest starting with the reworking of the existing implementation with StmtVisitor class and after that step-by-step extend the functionality of the visitor with handling other kinds а expressions.

clang/lib/Sema/SemaOpenMP.cpp

15274–15276

Again, this definitely should be reworked.

15293

The check is really bad. If you want to handle something like *(B+x) better to treat, say, B as the base and x as the index and, thus, treat the whole expression as something like B[x]

15303

Why do you need this?

15338

Do not call IgnoreParenImpCasts() here, just Visit(E). Also, I assume, it may lead to infinite recursion.

15356

Same here

[cchen]

Quoted Text

clang/lib/Sema/SemaOpenMP.cpp

15523

• (a ? b : c).
  • Error out for now, not sure what to for this one
• __builtin_choose_expr(a, b, c).
  • Allow this one since we know what to send to device at compile time
• a = b.
  • Sema accept this one, but need more work for codegen I think (Clang not emit the a=b ir if "a=b" inside map/motion clause)
• a?:b
  • What's the name of this kind of expression? I don't know what to do for this one
• __builtin_convertvector(x, ty)
  • not sure
• v.xy, where v is an extended vector
  • not sure
• a.b, where b is a bitfield
  • Error out for this one since bitfield is not addressable
• __builtin_bit_cast(v, ty)
  • Error out for this one since bitfield is not addressable
• casts
  • Error out if the expression inside the cast is not lvalue
• typeid() (also is an lvalue).
  • I guess we can support this just like normal variable?
• __uuidof(*comPtr)
  • not sure
• lambda calls
  • OpenMP 5.0 spec has rules for lambda. Now not error out for sema but need more work for codegen
• User defined literals
  • error out since not addressable

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15293

I don't understand how to do this, is there any sample code that I can learn from? Thanks

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15293

Not sure we have something like this, need to invent something new if we want to support lvalues in full. Or discard it as unsupported if we're unable to support it properly.

clang/include/clang/Basic/DiagnosticSemaKinds.td

def warn_omp_declare_variant_marked_as_declare_variant : Warning<
   "variant function in '#pragma omp declare variant' is itself marked as '#pragma omp declare variant'"
   >, InGroup<SourceUsesOpenMP>;
 def note_omp_marked_declare_variant_here : Note<"marked as 'declare variant' here">;
 def err_omp_one_defaultmap_each_category: Error<
   "at most one defaultmap clause for each variable-category can appear on the directive">;
 def err_omp_lastprivate_conditional_non_scalar : Error<
   "expected list item of scalar type in 'lastprivate' clause with 'conditional' modifier"
   >;
+def err_omp_non_lvalue_in_map_or_motion_clauses: Error<
+  "expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'"
+  >;
 } // end of OpenMP category
 
 let CategoryName = "Related Result Type Issue" in {
 // Objective-C related result type compatibility
 def warn_related_result_type_compatibility_class : Warning<
   "method is expected to return an instance of its class type "
   "%diff{$, but is declared to return $|"
   ", but is declared to return different type}0,1">;

clang/lib/CodeGen/CGOpenMPRuntime.cpp

void generateInfoForComponentList(
     //   float f[50];
     // }
     // struct S2 {
     //   int i;
     //   float f[50];
     //   S1 s;
     //   double *p;
     //   struct S2 *ps;
+    //   struct S2 **ps_ptr;
     // }
     // S2 s;
     // S2 *ps;
     //
     // map(d)
     // &d, &d, sizeof(double), TARGET_PARAM | TO | FROM
     //
     // map(i)
void generateInfoForComponentList(
     // &s, &(s.f[0]), 50*sizeof(float) + sizeof(struct S1), TARGET_PARAM
     // &s, &(s.f[0]), 22*sizeof(float), MEMBER_OF(1) | FROM
     // &s, &(s.s), sizeof(struct S1), MEMBER_OF(1) | FROM
     // ps, &(ps->p), sizeof(S2*), TARGET_PARAM
     // ps, &(ps->p), sizeof(double*), MEMBER_OF(4) (*)
     // &(ps->p), &(ps->p[0]), 33*sizeof(double), MEMBER_OF(4) | PTR_AND_OBJ | TO
     // (*) the struct this entry pertains to is the 4th element in the list
     //     of arguments, hence MEMBER_OF(4)
+    //
+    // map(*p)
+    // &p, &(p[0]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map(*(p+3))
+    // &p, &(p[3]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map((p+3)+7)
+    // &p, &(p[10]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map(*(ps->p))
+    // ps, &(ps->p[0]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map(*(ps->p+3))
+    // ps, &(ps->p[3]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map((ps->p+3)[1])
+    // ps, &(ps->p[4]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map(*(ps->ps->p))
+    // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM
+    // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1)
+    // &(ps->ps), &(ps->ps->p[0]), sizeof(double), MEMBER_OF(1) | PTR_AND_OBJ |
+    //  FROM
+    //
+    // map((ps->ps->p)[3])
+    // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM
+    // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1)
+    // &(ps->ps), &(ps->ps->p[3]), sizeof(double), MEMBER_OF(1) | PTR_AND_OBJ |
+    //  FROM
+    //
+    // map(*(ps->ps->p+3))
+    // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM
+    // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1)
+    // &(ps->ps), &(ps->ps->p[3]), sizeof(double), MEMBER_OF(1) | PTR_AND_OBJ |
+    //  FROM
+    //
+    // map(*(ps->ps->ps_ptr))
+    // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM
+    // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1)
+    // &(ps->ps_ptr), &(ps->ps->ps_ptr[0]), sizeof(S2*), MEMBER_OF(1) |
+    //   PTR_AND_OBJ | TO | FROM
 
     // Track if the map information being generated is the first for a capture.
     bool IsCaptureFirstInfo = IsFirstComponentList;
     // When the variable is on a declare target link or in a to clause with
     // unified memory, a reference is needed to hold the host/device address
     // of the variable.
     bool RequiresReference = false;
 
for (; I != CE; ++I) {
       bool IsPointer =
           (OASE && OMPArraySectionExpr::getBaseOriginalType(OASE)
                        .getCanonicalType()
                        ->isAnyPointerType()) ||
           I->getAssociatedExpression()->getType()->isAnyPointerType();
 
       if (Next == CE || IsPointer || IsFinalArraySection) {
         // If this is not the last component, we expect the pointer to be
         // associated with an array expression or member expression.
-        assert((Next == CE ||
-                isa<MemberExpr>(Next->getAssociatedExpression()) ||
-                isa<ArraySubscriptExpr>(Next->getAssociatedExpression()) ||
-                isa<OMPArraySectionExpr>(Next->getAssociatedExpression())) &&
+        assert((Next == CE || Next->getAssociatedExpression()->isLValue()) &&
                "Unexpected expression");
 
         Address LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
                          .getAddress(CGF);
 
         // If this component is a pointer inside the base struct then we don't
         // need to create any entry for it - it will be combined with the object
         // it is pointing to into a single PTR_AND_OBJ entry.
         bool IsMemberPointer =
for (; I != CE; ++I) {
 
           Types.push_back(Flags);
         }
 
         // If we have encountered a member expression so far, keep track of the
         // mapped member. If the parent is "*this", then the value declaration
         // is nullptr.
         if (EncounteredME) {
+          assert(!isa<FunctionDecl>(EncounteredME->getMemberDecl()) &&
+                 "Unexpected FunctionDecl");
           const auto *FD = dyn_cast<FieldDecl>(EncounteredME->getMemberDecl());
           unsigned FieldIndex = FD->getFieldIndex();
 
           // Update info about the lowest and highest elements for this struct
           if (!PartialStruct.Base.isValid()) {
             PartialStruct.LowestElem = {FieldIndex, LB};
             PartialStruct.HighestElem = {FieldIndex, LB};
             PartialStruct.Base = BP;
           } else if (FieldIndex < PartialStruct.LowestElem.first) {
             PartialStruct.LowestElem = {FieldIndex, LB};
           } else if (FieldIndex > PartialStruct.HighestElem.first) {
             PartialStruct.HighestElem = {FieldIndex, LB};
           }
         }
 
         // If we have a final array section, we are done with this expression.
         if (IsFinalArraySection)
           break;
 
         // The pointer becomes the base for the next element.

clang/lib/Sema/SemaOpenMP.cpp

static bool checkArrayExpressionDoesNotReferToUnitySize(Sema &SemaRef,
   Expr::EvalResult Result;
   if (!Length->EvaluateAsInt(Result, SemaRef.getASTContext()))
     return false; // Can't get the integer value as a constant.
 
   llvm::APSInt ConstLength = Result.Val.getInt();
   return ConstLength.getSExtValue() != 1;
 }
 
-// Return the expression of the base of the mappable expression or null if it
-// cannot be determined and do all the necessary checks to see if the expression
-// is valid as a standalone mappable expression. In the process, record all the
-// components of the expression.
-static const Expr *checkMapClauseExpressionBase(
-    Sema &SemaRef, Expr *E,
-    OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
-    OpenMPClauseKind CKind, bool NoDiagnose) {
-  SourceLocation ELoc = E->getExprLoc();
-  SourceRange ERange = E->getSourceRange();
-
-  // The base of elements of list in a map clause have to be either:
-  //  - a reference to variable or field.
-  //  - a member expression.
-  //  - an array expression.
-  //
-  // E.g. if we have the expression 'r.S.Arr[:12]', we want to retrieve the
-  // reference to 'r'.
-  //
-  // If we have:
-  //
-  // struct SS {
-  //   Bla S;
-  //   foo() {
-  //     #pragma omp target map (S.Arr[:12]);
-  //   }
-  // }
-  //
-  // We want to retrieve the member expression 'this->S';
+namespace {
 
-  const Expr *RelevantExpr = nullptr;
-
-  // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.2]
-  //  If a list item is an array section, it must specify contiguous storage.
-  //
-  // For this restriction it is sufficient that we make sure only references
-  // to variables or fields and array expressions, and that no array sections
-  // exist except in the rightmost expression (unless they cover the whole
-  // dimension of the array). E.g. these would be invalid:
-  //
-  //   r.ArrS[3:5].Arr[6:7]
-  //
-  //   r.ArrS[3:5].x
-  //
-  // but these would be valid:
-  //   r.ArrS[3].Arr[6:7]
-  //
-  //   r.ArrS[3].x
+class MapBaseChecker final : public StmtVisitor<MapBaseChecker, bool> {
+  Sema &SemaRef;
+  OpenMPClauseKind CKind;
+  OMPClauseMappableExprCommon::MappableExprComponentList &Components;
+  bool NoDiagnose;
+  const Expr *RelevantExpr{nullptr};
+  size_t Depth{0};
+  SmallVector<int, 4> CntStack;
+  bool AllowUnitySizeArraySection{true};
+  bool AllowWholeSizeArraySection{true};
+  SourceLocation ELoc;
+  SourceRange ERange;
 
-  bool AllowUnitySizeArraySection = true;
-  bool AllowWholeSizeArraySection = true;
-
-  while (!RelevantExpr) {
-    E = E->IgnoreParenImpCasts();
-
-    if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
-      if (!isa<VarDecl>(CurE->getDecl()))
-        return nullptr;
-
-      RelevantExpr = CurE;
+  void emitErrorMsg() {
+    if (!NoDiagnose) {
+      // If nothing else worked, this is not a valid map clause expression.
+      if (SemaRef.getLangOpts().OpenMP < 50) {
+        SemaRef.Diag(
+          ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
+          << ERange;
+      } else {
+        SemaRef.Diag(ELoc, diag::err_omp_non_lvalue_in_map_or_motion_clauses)
+          << ERange;
+      }
+    }
+  }
 
-      // If we got a reference to a declaration, we should not expect any array
-      // section before that.
-      AllowUnitySizeArraySection = false;
-      AllowWholeSizeArraySection = false;
+public:
+  bool VisitDeclRefExpr(DeclRefExpr *DRE) {
+    if (!isa<VarDecl>(DRE->getDecl()))
+      return false;
 
+    if (SemaRef.getLangOpts().OpenMP < 50) {
+      RelevantExpr = DRE;
       // Record the component.
-      CurComponents.emplace_back(CurE, CurE->getDecl());
-    } else if (auto *CurE = dyn_cast<MemberExpr>(E)) {
-      Expr *BaseE = CurE->getBase()->IgnoreParenImpCasts();
+      Components.emplace_back(DRE, DRE->getDecl());
+    } else {
+      const clang::Type *PT = DRE->getType().getTypePtr();
+      const std::string TypeStr = PT->getCanonicalTypeInternal().getAsString();
+      size_t TypeDecorCnt = std::count_if(TypeStr.begin(), TypeStr.end(),
+                                 [](char c) { return c == '*' || c == '['; });
+      // Normally we want to add Components if Depth == TypeDecorCnt, for example:
+      //
+      // int **B, *l;
+      // map(**(B + *l))
+      //   Depth for B: 2
+      //   TypeDecorCnt for B: 2
+      //   Depth for l: 3
+      //   TypeDecorCnt for l: 1
+      //
+      // However, combining with address_of operator(&), we might have Depth
+      // less than TypeDecorCnt as LValue
+      //
+      // int B;
+      // map(*(&B))
+      //  Depth for B: 1
+      //  TypeDecorCnt for B: 0
+      if (Depth <= TypeDecorCnt) {
+        RelevantExpr = cast<Expr>(DRE);
+        // record the component.
+        Components.emplace_back(DRE, DRE->getDecl());
+      }
+    }
+
+    return true;
+  }
+  bool VisitMemberExpr(MemberExpr *ME) {
+    Expr *E = cast<Expr>(ME);
+    Expr *BaseE = ME->getBase()->IgnoreParenImpCasts();
+    CntStack.emplace_back(Depth);
+    Depth = 0;
 
-      if (isa<CXXThisExpr>(BaseE))
-        // We found a base expression: this->Val.
-        RelevantExpr = CurE;
-      else
-        E = BaseE;
+      RelevantExpr = ME;
 
-      if (!isa<FieldDecl>(CurE->getMemberDecl())) {
-        if (!NoDiagnose) {
-          SemaRef.Diag(ELoc, diag::err_omp_expected_access_to_data_field)
-              << CurE->getSourceRange();
-          return nullptr;
-        }
-        if (RelevantExpr)
-          return nullptr;
-        continue;
-      }
+    if (!isa<FieldDecl>(ME->getMemberDecl())) {
+          << ME->getSourceRange();
+        return false;
+        return false;
+      return Visit(E->IgnoreParenImpCasts());;
 
-      auto *FD = cast<FieldDecl>(CurE->getMemberDecl());
+    auto *FD = cast<FieldDecl>(ME->getMemberDecl());
 
-      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.3]
-      //  A bit-field cannot appear in a map clause.
-      //
-      if (FD->isBitField()) {
-        if (!NoDiagnose) {
-          SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_clause)
-              << CurE->getSourceRange() << getOpenMPClauseName(CKind);
-          return nullptr;
-        }
-        if (RelevantExpr)
-          return nullptr;
-        continue;
-      }
+          << ME->getSourceRange() << getOpenMPClauseName(CKind);
+        return false;
+        return false;
+      return Visit(E->IgnoreParenImpCasts());;
 
-      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
-      //  If the type of a list item is a reference to a type T then the type
-      //  will be considered to be T for all purposes of this clause.
-      QualType CurType = BaseE->getType().getNonReferenceType();
 
-      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.2]
-      //  A list item cannot be a variable that is a member of a structure with
-      //  a union type.
-      //
-      if (CurType->isUnionType()) {
-        if (!NoDiagnose) {
-          SemaRef.Diag(ELoc, diag::err_omp_union_type_not_allowed)
-              << CurE->getSourceRange();
-          return nullptr;
-        }
-        continue;
-      }
+          << ME->getSourceRange();
+        return false;
+      return Visit(E->IgnoreParenImpCasts());;
 
-      // If we got a member expression, we should not expect any array section
-      // before that:
-      //
-      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.7]
-      //  If a list item is an element of a structure, only the rightmost symbol
-      //  of the variable reference can be an array section.
-      //
-      AllowUnitySizeArraySection = false;
-      AllowWholeSizeArraySection = false;
 
-      // Record the component.
-      CurComponents.emplace_back(CurE, FD);
-    } else if (auto *CurE = dyn_cast<ArraySubscriptExpr>(E)) {
-      E = CurE->getBase()->IgnoreParenImpCasts();
+    Components.emplace_back(ME, FD);
+    if (RelevantExpr || Visit(E->IgnoreParenImpCasts())) {
+      Depth = CntStack.back();
+      CntStack.pop_back();
+      return true;
+    } else {
+      return false;
+    }
+  }
+  bool VisitCXXThisExpr(CXXThisExpr *CTE) {
+    if (!RelevantExpr) {
+      Components.emplace_back(CTE, nullptr);
+      this->RelevantExpr = cast<Expr>(CTE);
+    }
+    return true;
+  }
+  bool VisitUnaryOperator(UnaryOperator *UO) {
+    if (SemaRef.getLangOpts().OpenMP < 50) {
+      emitErrorMsg();
+      return false;
+    }
+    if (UO->getOpcode() == UO_Deref || UO->getOpcode() == UO_AddrOf) {
+      if (!this->RelevantExpr && UO->getOpcode() == UO_Deref) {
+        Depth++;
+        // Record the component.
+        Components.emplace_back(UO, nullptr);
+      } else {
+        Depth--;
+      }
+      Expr *SubE = UO->getSubExpr()->IgnoreParenImpCasts();
+      if (SubE && Visit(SubE->IgnoreParenImpCasts())) {
+        if (UO->getOpcode() == UO_Deref)
+          Depth--;
+        else
+          Depth++;
+        return true;
+      }
+    }
+    return false;
+  }
+  bool VisitArraySubscriptExpr(ArraySubscriptExpr *AE) {
+    Depth++;
+    Expr *E = AE->getBase()->IgnoreParenImpCasts();
 
-      if (!E->getType()->isAnyPointerType() && !E->getType()->isArrayType()) {
-        if (!NoDiagnose) {
-          SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name)
-              << 0 << CurE->getSourceRange();
-          return nullptr;
-        }
-        continue;
-      }
+            << 0 << AE->getSourceRange();
+        return false;
+      return Visit(E->IgnoreParenImpCasts());
 
-      // If we got an array subscript that express the whole dimension we
-      // can have any array expressions before. If it only expressing part of
-      // the dimension, we can only have unitary-size array expressions.
-      if (checkArrayExpressionDoesNotReferToWholeSize(SemaRef, CurE,
-                                                      E->getType()))
-        AllowWholeSizeArraySection = false;
+    if (checkArrayExpressionDoesNotReferToWholeSize(SemaRef, AE,
 
-      if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
-        Expr::EvalResult Result;
-        if (CurE->getIdx()->EvaluateAsInt(Result, SemaRef.getASTContext())) {
-          if (!Result.Val.getInt().isNullValue()) {
-            SemaRef.Diag(CurE->getIdx()->getExprLoc(),
-                         diag::err_omp_invalid_map_this_expr);
-            SemaRef.Diag(CurE->getIdx()->getExprLoc(),
-                         diag::note_omp_invalid_subscript_on_this_ptr_map);
-          }
-        }
-        RelevantExpr = TE;
-      }
+      if (AE->getIdx()->EvaluateAsInt(Result, SemaRef.getASTContext())) {
+          SemaRef.Diag(AE->getIdx()->getExprLoc(),
+          SemaRef.Diag(AE->getIdx()->getExprLoc(),
 
-      // Record the component - we don't have any declaration associated.
-      CurComponents.emplace_back(CurE, nullptr);
-    } else if (auto *CurE = dyn_cast<OMPArraySectionExpr>(E)) {
-      assert(!NoDiagnose && "Array sections cannot be implicitly mapped.");
-      E = CurE->getBase()->IgnoreParenImpCasts();
+    Components.emplace_back(AE, nullptr);
+
+    if (E && Visit(E->IgnoreParenImpCasts())) {
+      Depth--;
+      return true;
+    }
+    return false;
+  }
+  bool VisitOMPArraySectionExpr(OMPArraySectionExpr *OASE) {
+    Depth++;
+    Expr *E = OASE->getBase()->IgnoreParenImpCasts();
 
-      QualType CurType =
-          OMPArraySectionExpr::getBaseOriginalType(E).getCanonicalType();
 
-      // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1]
-      //  If the type of a list item is a reference to a type T then the type
-      //  will be considered to be T for all purposes of this clause.
-      if (CurType->isReferenceType())
-        CurType = CurType->getPointeeType();
 
-      bool IsPointer = CurType->isAnyPointerType();
 
-      if (!IsPointer && !CurType->isArrayType()) {
-        SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name)
-            << 0 << CurE->getSourceRange();
-        return nullptr;
-      }
+        << 0 << OASE->getSourceRange();
+      return false;
 
-      bool NotWhole =
-          checkArrayExpressionDoesNotReferToWholeSize(SemaRef, CurE, CurType);
-      bool NotUnity =
-          checkArrayExpressionDoesNotReferToUnitySize(SemaRef, CurE, CurType);
+      checkArrayExpressionDoesNotReferToWholeSize(SemaRef, OASE, CurType);
+      checkArrayExpressionDoesNotReferToUnitySize(SemaRef, OASE, CurType);
 
-      if (AllowWholeSizeArraySection) {
-        // Any array section is currently allowed. Allowing a whole size array
-        // section implies allowing a unity array section as well.
-        //
-        // If this array section refers to the whole dimension we can still
-        // accept other array sections before this one, except if the base is a
-        // pointer. Otherwise, only unitary sections are accepted.
-        if (NotWhole || IsPointer)
-          AllowWholeSizeArraySection = false;
-      } else if (AllowUnitySizeArraySection && NotUnity) {
-        // A unity or whole array section is not allowed and that is not
-        // compatible with the properties of the current array section.
-        SemaRef.Diag(
-            ELoc, diag::err_array_section_does_not_specify_contiguous_storage)
-            << CurE->getSourceRange();
-        return nullptr;
-      }
+        << OASE->getSourceRange();
+      return false;
 
-      if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
-        Expr::EvalResult ResultR;
-        Expr::EvalResult ResultL;
-        if (CurE->getLength()->EvaluateAsInt(ResultR,
-                                             SemaRef.getASTContext())) {
-          if (!ResultR.Val.getInt().isOneValue()) {
-            SemaRef.Diag(CurE->getLength()->getExprLoc(),
-                         diag::err_omp_invalid_map_this_expr);
-            SemaRef.Diag(CurE->getLength()->getExprLoc(),
-                         diag::note_omp_invalid_length_on_this_ptr_mapping);
-          }
-        }
-        if (CurE->getLowerBound() && CurE->getLowerBound()->EvaluateAsInt(
-                                        ResultL, SemaRef.getASTContext())) {
-          if (!ResultL.Val.getInt().isNullValue()) {
-            SemaRef.Diag(CurE->getLowerBound()->getExprLoc(),
-                         diag::err_omp_invalid_map_this_expr);
-            SemaRef.Diag(CurE->getLowerBound()->getExprLoc(),
-                         diag::note_omp_invalid_lower_bound_on_this_ptr_mapping);
-          }
-        }
-        RelevantExpr = TE;
-      }
+      if (OASE->getLength()->EvaluateAsInt(ResultR,
+          SemaRef.Diag(OASE->getLength()->getExprLoc(),
+          SemaRef.Diag(OASE->getLength()->getExprLoc(),
+      if (OASE->getLowerBound() && OASE->getLowerBound()->EvaluateAsInt(
+          SemaRef.Diag(OASE->getLowerBound()->getExprLoc(),
+          SemaRef.Diag(OASE->getLowerBound()->getExprLoc(),
 
-      // Record the component - we don't have any declaration associated.
-      CurComponents.emplace_back(CurE, nullptr);
-    } else {
-      if (!NoDiagnose) {
-        // If nothing else worked, this is not a valid map clause expression.
-        SemaRef.Diag(
-            ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
-            << ERange;
-      }
-      return nullptr;
+    Components.emplace_back(OASE, nullptr);
+
+    if (RelevantExpr || Visit(E->IgnoreParenImpCasts())) {
+      Depth--;
+      return true;
+    return false;
+  }
+  bool VisitBinaryOperator(BinaryOperator *BO) {
+    if (SemaRef.getLangOpts().OpenMP < 50) {
+      emitErrorMsg();
+      return false;
     }
+    if (BO->getOpcode() == BO_Assign) {
+      this->RelevantExpr = cast<Expr>(BO);
+      // Record the component.
+      Components.emplace_back(BO, nullptr);
+      return true;
     }
-
+    Expr *LE = BO->getLHS()->IgnoreParenImpCasts();
+    Expr *RE = BO->getRHS()->IgnoreParenImpCasts();
+    if (!LE && !RE)
+      return false;
+    if (LE && RE) {
+      Visit(LE);
+      Visit(RE);
+      return true;
+    }
+    return false;
+  }
+  bool VisitConditionalOperator(ConditionalOperator *CO) {
+    emitErrorMsg();
+    return false;
+  }
+  bool VisitCallExpr(CallExpr *CE) {
+    emitErrorMsg();
+    return false;
+  }
+  bool VisitCastExpr(CastExpr *CE) {
+    if (SemaRef.getLangOpts().OpenMP < 50) {
+      emitErrorMsg();
+      return false;
+    }
+    if (!CE->isLValue())
+      return false;
+    return Visit(CE->getSubExpr());
+  }
+  bool VisitCXXMemberCallExpr(CXXMemberCallExpr *CCE) {
+    emitErrorMsg();
+    return false;
+  }
+  const Expr *getFoundBase() const {
     return RelevantExpr;
   }
+  OMPClauseMappableExprCommon::MappableExprComponentList getComponents() {
+    return Components;
+  }
+  explicit MapBaseChecker(
+      Sema &SemaRef, OpenMPClauseKind CKind,
+      OMPClauseMappableExprCommon::MappableExprComponentList &Components,
+      bool NoDiagnose, SourceLocation &ELoc, SourceRange &ERange)
+      : SemaRef(SemaRef), CKind(CKind), Components(Components),
+        NoDiagnose(NoDiagnose), ELoc(ELoc), ERange(ERange) {}
+};
+} // namespace
+
+// Return the expression of the base of the mappable expression or null if it
+// cannot be determined and do all the necessary checks to see if the expression
+// is valid as a standalone mappable expression. In the process, record all the
+// components of the expression.
+static const Expr *checkMapClauseExpressionBase(
+    Sema &SemaRef, Expr *E,
+    OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
+    OpenMPClauseKind CKind, bool NoDiagnose) {
+  SourceLocation ELoc = E->getExprLoc();
+  SourceRange ERange = E->getSourceRange();
+  MapBaseChecker Checker(SemaRef, CKind, CurComponents, NoDiagnose, ELoc,
+                         ERange);
+  if (Checker.Visit(E->IgnoreParenImpCasts())) {
+    return Checker.getFoundBase();
+  } else {
+    return nullptr;
+  }
+}
 
 // Return true if expression E associated with value VD has conflicts with other
 // map information.
 static bool checkMapConflicts(
     Sema &SemaRef, DSAStackTy *DSAS, const ValueDecl *VD, const Expr *E,
     bool CurrentRegionOnly,
     OMPClauseMappableExprCommon::MappableExprComponentListRef CurComponents,
     OpenMPClauseKind CKind) {
if (VE->isValueDependent() || VE->isTypeDependent() ||
       MVLI.UDMapperList.push_back(ER.get());
       // We can only analyze this information once the missing information is
       // resolved.
       MVLI.ProcessedVarList.push_back(RE);
       continue;
     }
 
     Expr *SimpleExpr = RE->IgnoreParenCasts();
+    Expr::isModifiableLvalueResult IsLV =
+        VE->isModifiableLvalue(SemaRef.Context, &ELoc);
 
-    if (!RE->IgnoreParenImpCasts()->isLValue()) {
-      SemaRef.Diag(ELoc,
-                   diag::err_omp_expected_named_var_member_or_array_expression)
+    if (!RE->IgnoreParenImpCasts()->isLValue() &&
+        IsLV != Expr::MLV_LValueCast) {
+      if (SemaRef.getLangOpts().OpenMP < 50) {
+        SemaRef.Diag(
+            ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
+            << RE->getSourceRange();
+      } else {
+        SemaRef.Diag(ELoc, diag::err_omp_non_lvalue_in_map_or_motion_clauses)
             << RE->getSourceRange();
+      }
       continue;
     }
 
     OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
     ValueDecl *CurDeclaration = nullptr;
 
     // Obtain the array or member expression bases if required. Also, fill the
     // components array with all the components identified in the process.
if (const auto *TE = dyn_cast<CXXThisExpr>(BE)) {
       // Skip restriction checking for variable or field declarations
       MVLI.ProcessedVarList.push_back(RE);
       MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
       MVLI.VarComponents.back().append(CurComponents.begin(),
                                        CurComponents.end());
       MVLI.VarBaseDeclarations.push_back(nullptr);
       continue;
     }
+    if (const auto *BO = dyn_cast<BinaryOperator>(BE)) {
+      MVLI.UDMapperList.push_back(nullptr);
+      MVLI.ProcessedVarList.push_back(RE);
+      MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
+      MVLI.VarComponents.back().append(CurComponents.begin(),
+                                       CurComponents.end());
+      MVLI.VarBaseDeclarations.push_back(nullptr);
+      continue;
+    }
 
     // For the following checks, we rely on the base declaration which is
     // expected to be associated with the last component. The declaration is
     // expected to be a variable or a field (if 'this' is being mapped).
     CurDeclaration = CurComponents.back().getAssociatedDeclaration();
     assert(CurDeclaration && "Null decl on map clause.");
     assert(
         CurDeclaration->isCanonicalDecl() &&
         "Expecting components to have associated only canonical declarations.");
 
     auto *VD = dyn_cast<VarDecl>(CurDeclaration);
     const auto *FD = dyn_cast<FieldDecl>(CurDeclaration);
 
     assert((VD || FD) && "Only variables or fields are expected here!");
     (void)FD;
 
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.10]
     // threadprivate variables cannot appear in a map clause.
     // OpenMP 4.5 [2.10.5, target update Construct]
     // threadprivate variables cannot appear in a from clause.
     if (VD && DSAS->isThreadPrivate(VD)) {

clang/test/OpenMP/target_map_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
 // RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-mapping -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) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    #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 lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}}
     {}
-    #pragma omp target map(arg,(c+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    #pragma omp target map(arg,(c+1)[2]) // le45-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 SAclient(int arg) {
   #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]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  #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 lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}}
   {}
   #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.
void SAclient(int arg) {
   #pragma omp target map(r.ArrS[3].Arr[6:7])
   {}
   #pragma omp target map(r.S.Arr[4:5])
   {}
   #pragma omp target map(r.S.Ptr[4:5])
   {}
   #pragma omp target map(r.S.Ptr[:])  // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
   {}
-  #pragma omp target map((p+1)->A)  // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  #pragma omp target map((p+1)->A)  // le45-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();
   }

 #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) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#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 lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}}
   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) // expected-error 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) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}}
 #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(S2::S2sc)
#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]) // expected-error {{xpected 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]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}}
 #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)

clang/test/OpenMP/target_teams_map_messages.cpp

 // RUN: %clang_cc1 -verify=expected,le45 -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 -fopenmp-version=50 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp -Wuninitialized
 
 // RUN: %clang_cc1 -verify=expected,le45 -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,le45 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-mapping -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 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) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    #pragma omp target teams map(arg,a*2) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
     {}
-    #pragma omp target teams map(arg,(c+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    #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}}
     {}
     #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]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  #pragma omp target teams map(r.S.foo()[:12]) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
   #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)  // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  #pragma omp target teams map((p+1)->A) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
   {}
   #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) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams map(I) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
   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) // expected-error 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) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+
 #pragma omp target data map(argc)
 #pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
 #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
 #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(S2::S2sc)

 #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]) // expected-error {{xpected 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]) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target data map(argc)
 #pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
 #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
 #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)

clang/test/OpenMP/target_update_codegen.cpp

void device_side_scan(int arg) {
   // CK4: tgt_target_data_update
   // CK4: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
   // CK4: ret
   // TCK4-NOT: tgt_target_data_update
   #pragma omp target update from(arg) if(arg) device(4)
   {++arg;}
 }
 #endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-64
+// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-64
+// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32
+// RUN: %clang_cc1 -DCK5 -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 CK5 --check-prefix CK5-32
+
+// RUN: %clang_cc1 -DCK5 -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-ONLY0 %s
+// RUN: %clang_cc1 -DCK5 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -fopenmp-version=50 -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 -DCK5 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK5
+
+// CK5: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+// CK5-LABEL: lvalue
+void lvalue(int *B, int l, int e) {
+
+  // CK5-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK5-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK5-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK5-DAG: store i32* [[ZERO:%.+]], i32** [[BPC0]]
+  // CK5-DAG: store i32* [[ONE:%.+]], i32** [[PC0]]
+  // CK5-DAG: [[ZERO]] = load i32*, i32** [[B_ADDR:%.+]]
+  // CK5-DAG: [[ONE]] = load i32*, i32** [[B_ADDR]]
+  #pragma omp target update to(*B)
+  *B += e;
+  #pragma omp target update from(*B)
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK6 -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 CK6 --check-prefix CK6-64
+// RUN: %clang_cc1 -DCK6 -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 CK6 --check-prefix CK6-64
+// RUN: %clang_cc1 -DCK6 -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 CK6 --check-prefix CK6-32
+// RUN: %clang_cc1 -DCK6 -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 CK6 --check-prefix CK6-32
+
+// RUN: %clang_cc1 -DCK6 -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-ONLY0 %s
+// RUN: %clang_cc1 -DCK6 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK6 -verify -fopenmp-simd -fopenmp-version=50 -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 -DCK6 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK6
+
+// CK6: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+// CK6-LABEL: lvalue
+void lvalue(int *B, int l, int e) {
+
+  // CK6-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK6-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK6-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK6-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK6-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK6-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK6-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK6-DAG: store i32* [[ZERO:%.+]], i32** [[BPC0]]
+  // CK6-DAG: store i32* [[ADD_PTR:%.+]], i32** [[PC0]]
+  // CK6-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i{{32|64}} [[IDX_EXT:%.+]]
+  // CK6-DAG: [[TWO:%.+]] = load i32, i32* [[L_ADDR:%.+]]
+  // CK6-DAG: store i32 [[L:%.+]], i32* [[L_ADDR]]
+  #pragma omp target update to(*(B+l))
+  *(B+l) += e;
+  #pragma omp target update from(*(B+l))
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK7 -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 CK7 --check-prefix CK7-64
+// RUN: %clang_cc1 -DCK7 -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 CK7 --check-prefix CK7-64
+// RUN: %clang_cc1 -DCK7 -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 CK7 --check-prefix CK7-32
+// RUN: %clang_cc1 -DCK7 -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 CK7 --check-prefix CK7-32
+
+// RUN: %clang_cc1 -DCK7 -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-ONLY0 %s
+// RUN: %clang_cc1 -DCK7 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK7 -verify -fopenmp-simd -fopenmp-version=50 -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 -DCK7 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK7
+
+// CK7: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+// CK7-LABEL: lvalue
+void lvalue(int *B, int l, int e) {
+
+  // CK7-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK7-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK7-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK7-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK7-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK7-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK7-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK7-DAG: store i32* [[ZERO:%.+]], i32** [[BPC0]]
+  // CK7-DAG: store i32* [[ARRAY_IDX:%.+]], i32** [[PC0]]
+  // CK7-DAG: [[ZERO]] = load i32*, i32** [[B_ADDR:%.+]]
+  // CK7-DAG: [[ARRAY_IDX]] = getelementptr inbounds i32, i32* [[ADD_PTR:%.+]], i{{32|64}} [[IDX_PROM:%.+]]
+  // CK7-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i64 [[IDX_EXT:%.+]]
+  // CK7-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i32 [[TWO:%.+]] 
+  // CK7-32-DAG: [[ONE]] = load i32*, i32** [[B_ADDR]]
+  // CK7-32-DAG: [[TWO]] = load i32, i32* [[L_ADDR:%.+]]
+  // CK7-32-DAG: [[IDX_PROM]] = load i32, i32* [[L_ADDR]]
+  // CK7-64-DAG: [[IDX_EXT:%.+]] = sext i32 [[TWO:%.+]] to i64 
+  // CK7-64-DAG: [[TWO:%.+]] = load i32, i32* [[L_ADDR:%.+]]
+  // CK7-64-DAG: [[IDX_PROM]] = sext i32 [[THREE:%.+]] to i64
+  // CK7-64-DAG: [[THREE]] = load i32, i32* [[L_ADDR]]
+  #pragma omp target update to((B+l)[l])
+  (B+l)[l] += e;
+  #pragma omp target update from((B+l)[l])
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK8 -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 CK8 --check-prefix CK8-64
+// RUN: %clang_cc1 -DCK8 -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 CK8 --check-prefix CK8-64
+// RUN: %clang_cc1 -DCK8 -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 CK8 --check-prefix CK8-32
+// RUN: %clang_cc1 -DCK8 -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 CK8 --check-prefix CK8-32
+
+// RUN: %clang_cc1 -DCK8 -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-ONLY0 %s
+// RUN: %clang_cc1 -DCK8 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp-simd -fopenmp-version=50 -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 -DCK8 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK8
+// CK8-64: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} 8, i64 4]
+// CK8-32: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} 4, i64 4]
+// CK8: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17]
+
+// CK8-LABEL: lvalue
+void lvalue(int **B, int l, int e) {
+
+  // CK8-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}], [2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK8-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK8-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK8-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK8-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK8-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
+  // CK8-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK8-DAG: store i32** [[ARRAY_IDX_1:%.+]], i32*** [[BPC0]]
+  // CK8-DAG: store i32* [[ARRAY_IDX_4:%.+]], i32** [[PC0]]
+  // CK8-DAG: store i32** [[ARRAY_IDX_1]], i32*** [[NINE:%.+]]
+  // CK8-DAG: [[ARRAY_IDX_1]] = getelementptr inbounds i32*, i32** [[ADD_PTR:%.+]], i{{.+}} 1
+  // CK8-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i{{.+}} [[IDX_EXT:%.+]]
+  // CK8-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
+  // CK8-DAG: [[ARRAY_IDX_4]] = getelementptr inbounds i32, i32* [[FIVE:%.+]], i{{.+}} 2
+  // CK8-64-DAG: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds i32*, i32** [[ADD_PTR_2:%.+]], i{{.+}} 1
+  // CK8-64-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32*, i32** %3, i{{.+}} [[IDX_EXT_1:%.+]]
+  // CK8-64-DAG: [[IDX_EXT_1]] = sext i32 [[L_VAL:%.+]] to i{{.+}}
+  // CK8-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i{{.+}} [[L_VAL:%.+]]
+  // CK8-32-DAG: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds i32*, i32** [[ADD_PTR_2:%.+]], i{{.+}} 1
+  // CK8-32-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32*, i32** %3, i{{.+}} [[L_VAL:%.+]]
+  #pragma omp target update to((B+l)[1][2])
+  (B+l)[1][2] += e;
+  #pragma omp target update from((B+l)[1][2])
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK9 -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 CK9 --check-prefix CK9-64
+// RUN: %clang_cc1 -DCK9 -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 CK9 --check-prefix CK9-64
+// RUN: %clang_cc1 -DCK9 -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 CK9 --check-prefix CK9-32
+// RUN: %clang_cc1 -DCK9 -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 CK9 --check-prefix CK9-32
+
+// RUN: %clang_cc1 -DCK9 -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-ONLY0 %s
+// RUN: %clang_cc1 -DCK9 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK9 -verify -fopenmp-simd -fopenmp-version=50 -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 -DCK9 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK9
+
+struct S {
+  double *p;
+};
+
+// CK9: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} 32, i64 281474976710673]
+// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+
+// CK9-LABEL: lvalue
+void lvalue(struct S *s, int l, int e) {
+
+  // CK9-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK9-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK9-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK9-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+  //
+  // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK9-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK9-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
+  // CK9-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
+  // CK9-DAG: store double** [[P:%.+]], double*** [[BPC0]]
+  // CK9-DAG: store double* [[THREE:%.+]], double** [[PC0]]
+  // CK9-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE0]]
+  // CK9-DAG: [[P]] = getelementptr inbounds [[STRUCT_S:%.+]], [[STRUCT_S]]* [[ONE:%.+]], i32 0, i32 0
+  // CK9-DAG: [[ONE]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S_ADDR:%.+]]
+  // CK9-DAG: [[THREE]] = load double*, double** [[P_1:%.+]],
+  // CK9-DAG: [[P_1]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[TWO:%.+]], i32 0, i32 0
+  // CK9-DAG: [[TWO]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S_ADDR:%.+]]
+  #pragma omp target update to(*(s->p))
+    *(s->p) += e;
+  #pragma omp target update from(*(s->p))
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK10 -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 CK10 --check-prefix CK10-64
+// RUN: %clang_cc1 -DCK10 -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 CK10 --check-prefix CK10-64
+// RUN: %clang_cc1 -DCK10 -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 CK10 --check-prefix CK10-32
+// RUN: %clang_cc1 -DCK10 -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 CK10 --check-prefix CK10-32
+
+// RUN: %clang_cc1 -DCK10 -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-ONLY0 %s
+// RUN: %clang_cc1 -DCK10 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK10 -verify -fopenmp-simd -fopenmp-version=50 -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 -DCK10 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK10
+
+struct S {
+  double *p;
+};
+
+// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+
+// CK10-LABEL: lvalue
+void lvalue(struct S *s, int l, int e) {
+
+  // CK10-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK10-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK10-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK10-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+  //
+  // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK10-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK10-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
+  // CK10-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
+  // CK10-DAG: store double** [[P_VAL:%.+]], double*** [[BPC0]]
+  // CK10-DAG: store double* [[ADD_PTR:%.+]], double** [[PC0]]
+  // CK10-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE0]]
+  // CK10-64-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[IDX_EXT:%.+]]
+  // CK10-32-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[L_VAL:%.+]]
+  // CK10-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
+  // CK10-DAG: [[THREE]] = load double*, double** [[P_VAL_1:%.+]]
+  // CK10-DAG: getelementptr inbounds {{.+}}, {{.+}}* [[SS_1:%.+]], i32 0, i32 0
+  // CK10-DAG: [[SS_1]] = load {{.+}}*, {{.+}}** [[S_ADDR:%.+]]
+  #pragma omp target update to(*(s->p+l))
+    *(s->p+l) += e;
+  #pragma omp target update from(*(s->p+l))
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK11 -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 CK11 --check-prefix CK11-64
+// RUN: %clang_cc1 -DCK11 -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 CK11 --check-prefix CK11-64
+// RUN: %clang_cc1 -DCK11 -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 CK11 --check-prefix CK11-32
+// RUN: %clang_cc1 -DCK11 -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 CK11 --check-prefix CK11-32
+
+// RUN: %clang_cc1 -DCK11 -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-ONLY0 %s
+// RUN: %clang_cc1 -DCK11 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK11 -verify -fopenmp-simd -fopenmp-version=50 -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 -DCK11 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK11
+
+struct S {
+  double *p;
+};
+// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+
+// CK11-LABEL: lvalue
+void lvalue(struct S *s, int l, int e) {
+
+  // CK11-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK11-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK11-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK11-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+  //
+  // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK11-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK11-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
+  // CK11-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
+  // CK11-DAG: store double** [[P:%.+]], double*** [[BPC0]]
+  // CK11-DAG: store double* [[ARRAY_IDX:%.+]], double** [[PC0]]
+  // CK11-DAG: store i{{.+}} 8, i{{.+}} [[SIZE0]]
+  // CK11-DAG: [[P]] = getelementptr inbounds [[STRUCT_S:%.+]], [[STRUCT_S]]* [[SS_1:%.+]], i32 0, i32 0
+  // CK11-DAG: [[ARRAY_IDX]] = getelementptr inbounds double, double* [[ADD_PTR:%.+]], i{{.+}} 3
+  // CK11-64-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[IDX_EXT:%.+]]
+  // CK11-32-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[L_VAL:%.+]]
+  // CK11-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
+  #pragma omp target update to((s->p+l)[3])
+    (s->p+l)[3] += e;
+  #pragma omp target update from((s->p+l)[3])
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK12 -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 CK12 --check-prefix CK12-64
+// RUN: %clang_cc1 -DCK12 -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 CK12 --check-prefix CK12-64
+// RUN: %clang_cc1 -DCK12 -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 CK12 --check-prefix CK12-32
+// RUN: %clang_cc1 -DCK12 -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 CK12 --check-prefix CK12-32
+
+// RUN: %clang_cc1 -DCK12 -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-ONLY0 %s
+// RUN: %clang_cc1 -DCK12 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -DCK12 -verify -fopenmp-simd -fopenmp-version=50 -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 -DCK12 -fopenmp-simd -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-simd -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 --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK12
+
+struct S {
+  double *p;
+  struct S *sp;
+};
+// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710672, i64 17]
+
+// CK12-LABEL: lvalue
+void lvalue(struct S *s, int l, int e) {
+
+  // CK12-DAG: call void @__tgt_target_data_update(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK12-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK12-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK12-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+  //
+  // CK12-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+  // CK12-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+  // CK12-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+  // CK12-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to double***
+  // CK12-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double**
+  // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[BPC2]]
+  // CK12-DAG: store double* [[SIX:%.+]], double** [[PC2]]
+  // CK12-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE2]]
+  // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK12-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK12-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to [[STRUCT_S:%.+]]***
+  // CK12-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double***
+  // CK12-DAG: store [[STRUCT_S]]** [[SP:%.+]], [[STRUCT_S]]*** [[BPC1]]
+  // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[PC1]]
+  // CK12-DAG: store i{{.+}} {{4|8}}, i{{.+}}* [[SIZE1]]
+  // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK12-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0
+  // CK12-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[STRUCT_S:%.+]]**
+  // CK12-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to [[STRUCT_S]]***
+  // CK12-DAG: store [[STRUCT_S]]** [[SP:%.+]], [[STRUCT_S]]*** [[S_VAL:%.+]]
+  // CK12-DAG: store i{{.+}} {{.+}}, i{{.+}}* [[SIZE0]]
+  // CK12-DAG: [[SP]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[ONE:%.+]], i32 0, i32 1
+  #pragma omp target update to(*(s->sp->p))
+    *(s->sp->p) = e;
+  #pragma omp target update from(*(s->sp->p))
+}
+
+#endif
 #endif

clang/test/OpenMP/target_update_from_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 struct S7 {
   int i;
   int a[50];
   float x;
   S6 s6[5];
   double *p;
   unsigned bfa : 4;
 };
+struct S8 {
+  int *ptr;
+  int a;
+  struct S7* S;
+  void foo() {
+#pragma omp target update from(*this) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    {}
+#pragma omp target update from(*(&(*this))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(&(*(this->S->i+this->S->p)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(&(*(this->S->i+this->S->s6[0].pp)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(&(*(a+this->ptr)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(&(*(*(this->ptr)+a+this->ptr)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    {}
+  }
+};
+
+struct Base {
+  virtual ~Base() {}
+};
+struct Derived: Base {
+  virtual void name() {}
+};
+
 
 S3 h;
 #pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
 
 typedef int to;
 
 template <typename T, int I> // expected-note {{declared here}}
 T tmain(T argc) {

 #pragma omp target update from // expected-error {{expected '(' after 'from'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(x)
 #pragma omp target update from(t[:I])
 #pragma omp target update from(T) // expected-error {{'T' does not refer to a value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
-#pragma omp target update from(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(I) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target update from(S2::S2s)
 #pragma omp target update from(S2::S2sc)
 #pragma omp target update from(from)
 #pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}}
-#pragma omp target update from(argc > 0 ? x : y) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} 
+#pragma omp target update from(argc > 0 ? x : y) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+
 #pragma omp target update from(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
 #pragma omp target update from(ba)
 #pragma omp target update from(h) // expected-error {{threadprivate variables are not allowed in 'from' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(k), to(k) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
 #pragma omp target update from(t), from(t[:5]) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
 #pragma omp target update from(da)
 #pragma omp target update from(da[:4])
 
 #pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
 #pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
-#pragma omp target update from(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(x, (m+1)[2]) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target update from(s7.i, s7.a[:3])
 #pragma omp target update from(s7.s6[1].aa[0:5])
 #pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update from(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update from(s7.p[:10])
 #pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}}
 #pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
 #pragma omp target data map(to: s7.i)
   {
 #pragma omp target update from(s7.x)
   }
 
+  const int &rci = i;
+  Base *b = new Base;
+  int *pd = const_cast<int*>(&d);
+#pragma omp target update to(const_cast<int&>(rci))
+#pragma omp target update to(const_cast<int*>(pd))
+#pragma omp target update to(static_cast<double>(i))
+#pragma omp target update to(dynamic_cast<Derived*>(b))
+#pragma omp target update to(reinterpret_cast<int*>(b))
+  {}
+
   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, t[20];
   int &j = i;
   int *k = &j;
   int x;
   int y;
   int from;
   const int (&l)[5] = da;
   int *m;
   S7 s7;
+  int **BB, *offset;
+  int ***BBB;
 
 #pragma omp target update from // expected-error {{expected '(' after 'from'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(x)
 #pragma omp target update from(t[:i])
 #pragma omp target update from(S2::S2s)
 #pragma omp target update from(S2::S2sc)
 #pragma omp target update from(from)
 #pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}}
-#pragma omp target update from(argc > 0 ? x : y) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(argc > 0 ? x : y) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
 #pragma omp target update from(ba)
 #pragma omp target update from(h) // expected-error {{threadprivate variables are not allowed in 'from' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(k), to(k) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
 #pragma omp target update from(t), from(t[:5]) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
 #pragma omp target update from(da)
 #pragma omp target update from(da[:4])
 
 #pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
 #pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
-#pragma omp target update from(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update from(x, (m+1)[2]) // // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target update from(s7.i, s7.a[:3])
 #pragma omp target update from(s7.s6[1].aa[0:5])
 #pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update from(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update from(s7.p[:10])
 #pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}}
 #pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+#pragma omp target update from(*(&y)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(&(*(&y)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(**(BB+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(**(BB+y)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(m+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(**(*offset+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(**(y+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(*offset+m)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(**(*offset+BB+*m)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(**(*(*(&offset))+BB+*m)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(x+*(y+*(**BB+BBB)+s7.i)+static_cast<int>(*(s7.s6[0].pp)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target data map(to: s7.i)
   {
 #pragma omp target update from(s7.x)
   }
 
+  const int &rci = i;
+  Base *b = new Base;
+  int *pd = const_cast<int*>(&d);
+#pragma omp target update from(const_cast<int&>(rci))
+#pragma omp target update from(const_cast<int*>(pd))
+#pragma omp target update from(static_cast<double>(i))
+#pragma omp target update from(dynamic_cast<Derived*>(b))
+#pragma omp target update from(reinterpret_cast<int*>(b))
+  {}
+
+#pragma omp target update from(static_cast<double>(i + 1)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}}
+  {}
+
   return tmain<int, 3>(argc)+tmain<to, 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_update_to_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
-
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=40 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 %s -Wno-openmp-mapping -Wuninitialized
 
 void foo() {
 }
 
 bool foobool(int argc) {
   return argc;
 }
 

 struct S7 {
   int i;
   int a[50];
   float x;
   S6 s6[5];
   double *p;
   unsigned bfa : 4;
 };
+struct S8 {
+  int *ptr;
+  int a;
+  struct S7* S;
+  void foo() {
+#pragma omp target update to(*this) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    {}
+#pragma omp target update to(*(&(*this))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(&(*(this->S->i+this->S->p)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(&(*(this->S->i+this->S->s6[0].pp)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(&(*(a+this->ptr)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(&(*(*(this->ptr)+a+this->ptr)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    {}
+  }
+};
+
+struct Base {
+  virtual ~Base() {}
+};
+struct Derived: Base {
+  virtual void name() {}
+};
 
 S3 h;
 #pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
 
 typedef int from;
 
 template <typename T, int I> // expected-note {{declared here}}
 T tmain(T argc) {

 #pragma omp target update to // expected-error {{expected '(' after 'to'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(x)
 #pragma omp target update to(t[:I])
 #pragma omp target update to(T) // expected-error {{'T' does not refer to a value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
-#pragma omp target update to(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(I) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target update to(S2::S2s)
 #pragma omp target update to(S2::S2sc)
 #pragma omp target update to(to)
 #pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}}
-#pragma omp target update to(argc > 0 ? x : y) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} 
+#pragma omp target update to(argc > 0 ? x : y) // le50-error 2 {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target update to(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
 #pragma omp target update to(ba)
 #pragma omp target update to(h) // expected-error {{threadprivate variables are not allowed in 'to' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(k), from(k) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
 #pragma omp target update to(t), to(t[:5]) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}}
 #pragma omp target update to(da)
 #pragma omp target update to(da[:4])
 
 #pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
 #pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
-#pragma omp target update to(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(x, (m+1)[2]) // le45-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target update to(s7.i, s7.a[:3])
 #pragma omp target update to(s7.s6[1].aa[0:5])
 #pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update to(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update to(s7.p[:10])
 #pragma omp target update to(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'to' clause}}
 #pragma omp target update to(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
 #pragma omp target data map(to: s7.i)
   {
 #pragma omp target update to(s7.x)
   }
+
+  const int &rci = i;
+  Base *b = new Base;
+  int *pd = const_cast<int*>(&d);
+#pragma omp target update to(const_cast<int&>(rci))
+#pragma omp target update to(const_cast<int*>(pd))
+#pragma omp target update to(static_cast<double>(i))
+#pragma omp target update to(dynamic_cast<Derived*>(b))
+#pragma omp target update to(reinterpret_cast<int*>(b))
+  {}
   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, t[20];
   int &j = i;
   int *k = &j;
   int x;
   int y;
   int to;
   const int (&l)[5] = da;
   S7 s7;
   int *m;
+  int **BB, *offset;
+  int ***BBB;
 
 #pragma omp target update to // expected-error {{expected '(' after 'to'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(x)
 #pragma omp target update to(t[:i])
 #pragma omp target update to(S2::S2s)
 #pragma omp target update to(S2::S2sc)
 #pragma omp target update to(to)
 #pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}}
-#pragma omp target update to(argc > 0 ? x : y) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(argc > 0 ? x : y) // le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}} le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}}
 #pragma omp target update to(ba)
 #pragma omp target update to(h) // expected-error {{threadprivate variables are not allowed in 'to' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(k), from(k) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
 #pragma omp target update to(t), to(t[:5]) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}}
 #pragma omp target update to(da)
 #pragma omp target update to(da[:4])
 
 #pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}}
 #pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}}
-#pragma omp target update to(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target update to(x, (m+1)[2]) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
 #pragma omp target update to(s7.i, s7.a[:3])
 #pragma omp target update to(s7.s6[1].aa[0:5])
 #pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update to(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update to(s7.p[:10])
 #pragma omp target update to(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'to' clause}}
 #pragma omp target update to(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
 #pragma omp target data map(to: s7.i)
   {
 #pragma omp target update to(s7.x)
   }
 
+#pragma omp target update to(*(&y)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(&(*(&y)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(**(BB+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(**(BB+y)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(m+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(**(*offset+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(**(y+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(*offset+m)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(**(*offset+BB+*m)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(**(*(*(&offset))+BB+*m)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(x+*(y+*(**BB+BBB)+s7.i)+static_cast<int>(*(s7.s6[0].pp)))) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {}
+
+  const int &rci = i;
+  Base *b = new Base;
+  int *pd = const_cast<int*>(&d);
+#pragma omp target update to(const_cast<int&>(rci))
+#pragma omp target update to(const_cast<int*>(pd))
+#pragma omp target update to(static_cast<double>(i))
+#pragma omp target update to(dynamic_cast<Derived*>(b))
+#pragma omp target update to(reinterpret_cast<int*>(b))
+  {}
+
+#pragma omp target update to(static_cast<double>(i + 1)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} le50-error {{expected lvalue with no function call in '#pragma omp target update' and '#pragma omp target map'}}
+  {}
+
   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}}
 }