[OpenMP5.0] Allow pointer arithmetic in motion/map clause

[cchen]

Summary

Base declaration in pointer arithmetic expression is determined by
binary search with type information. Take "int *a, *b; *(a+*b)" as an
example, we determine the base by checking the type of LHS and RHS. In
this case the type of LHS is "int *", the type of RHS is "int",
therefore, we know that we need to visit LHS in order to find base
declaration.

[cchen]

Remove redundant code in test.

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7617

What about binary operator?

clang/lib/Sema/SemaOpenMP.cpp

15703

Better to discuss the implementation for this in a separate patch, it really requires some additional analysis and extra work.

15717

Again, bad idea to count this stuff.

15734

No need for this->

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7617

Didn't add check here since I didn't add any binary operator into Component in the MapBaseChecker class. Should I add it?

clang/lib/Sema/SemaOpenMP.cpp

15734

Will fix this.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15703

Why should we do this in another patch? The only thing this patch does is extending for pointer arithmetic (doing analysis on BinOp).

15717

Will it be better if just check if the subtree is an offset? So that we only need to check if it does not have any decorator in type?

[ABataev]

clang/include/clang/Basic/DiagnosticSemaKinds.td

10116

I would say that this lvalue must be addressable, no? Also, some function calls can be handled, probably, those returning rvalues, for example, or constexprs.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7617

I think so. Otherwise those binary expressions are meaningless. Or treated in some unexpected way.

clang/lib/Sema/SemaOpenMP.cpp

15703

You're doing some checks in this patch using counting of '*' and '[' in this patch. This is definitely not correct. That's why I suggest to move it in a separate patch and discus the solution there and do not block all other kinds of expressions.

15717

What do you mean? Exclude, say, RHS or LHS from the analysis, if it is a complex expression and you can use another part as a base? That's worth trying, at least.

[cchen]

Remove the counting in MapBaseChecker, instead, comparing the type of root with
the type of LHS and RHS, and only visit the subtree having the same type as root.
This scheme pass all the test I have. I'll add BinOp into components in the next
revision. Thanks

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15717

Generally yes, but the current implementation was the opposite way: visit the subtree with the same type as root (binop), if found, then we just visit it.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15688

I think you need to extend this to support paren expressions, at least.

15713–15715

I don't think you need these variables here anymore.

15719–15723

This does not look correct. At least, it should return the result of one of Visit() calls. Plus, it seems to me, it won't handle something like *(a+5). That's why I suggest excluding binary op from this patch, it is not complete.

[cchen]

Add BinaryOperator in components and fix logic in MapBaseChecker.

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7634

I'm now not checking for isLValue here since the binary expression without unary operator is not an LValue.

clang/lib/Sema/SemaOpenMP.cpp

15719–15723

Now return the result of Visit() calls. Also, this patch can handle *(a+5) or more complicated pointer arithmetic expression, such as CK6 and CK13 in target_update_codegen.cpp.

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15688

I have UO->getSubExpr()->IgnoreParenImpCasts() here, or I might not understand what you mean by "support paren expression".

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15688

(*(..)) is not supported, I think. Or, maybe, *((a)+b)

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15688

Both cases are handled correctly. For (*((b)+a)), the binop in VisitBinaryOperator()does not include the outer paren and for *((a)+b), the parens can be handled by Expr *SubE = UO->getSubExpr()->IgnoreParenImpCasts();.

[ABataev]

clang/include/clang/Basic/DiagnosticSemaKinds.td

10116

Fix this message, please, it is not quite correct.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7617–7634

Do you have the tests for these changes?

7750

Change this line to just cast<FieldDecl>. In this case the assert is not required.

clang/lib/Sema/SemaOpenMP.cpp

15688

I think we need to change the way we check the expression in call for MapBaseChecker::Visit in checkMapClauseExpressionBase or extend the check in the checkMapClauseExpressionBase. E->IgnoreParenImpCasts() may drop implicit lvalue-to-rvalue cast and we may allow mapping of rvalues.

15700–15701

I think, here we should have return RelevantExpr || Visit(UO->getSubExpr()) without dropping parens and implicit casts. Again, we may drop some important implicit casts

15708–15711

Add a check here that the type of the expression must be a pointer type.

15720–15722

Just return Visit(RE); without else

15723

Unreachable code

[cchen]

clang/include/clang/Basic/DiagnosticSemaKinds.td

10116

Will fix this, thanks.

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7617–7634

I think target_update_codegen.cpp has already tested those.

[ABataev]

clang/test/OpenMP/target_update_codegen.cpp

371–372

I suppose you missed some checks here, where IDX_EXT and TWO are used.

[cchen]

Fix based on feedback

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15695

Hmm, addr_of creates rvalues.

15698

No need for else here

15715

Not sure this should be an assert. What if you have something like *((a+b) + c) and a+b is not pointer arithmetic?

15725

No need for else here

clang/test/OpenMP/target_update_codegen.cpp

374

Must be just [[TWO]]?

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15715

Good point, I'll fix it. Thanks

clang/test/OpenMP/target_update_codegen.cpp

374

You mean use a more representable name?

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15715

Well, for this case, the type of binop expr *((a+b) + c) is int *, and the type of LHS is int, while the type of RHS is int *, therefore, expr (a+b) is not going to be visited. I'm using assert here since before the checkMapClauseExpressionBase be called, we have already check for whether the expression is lvalue. Maybe I'll just error out if binop here is not a pointer type instead of assert?

[ABataev]

clang/test/OpenMP/target_update_codegen.cpp

374

No, I mean that here you should not redefine the variable but check that the previously defined TWO is used in the check line.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15715

Yes

[cchen]

Fix test and add check for unary operator.

[cchen]

Remove redundant else

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15689–15690

UO_Deref is always r-value. So, !UO->isLValue() will be always triggered for UO_Deref. You can just exclude it from the condition.

clang/test/OpenMP/target_update_codegen.cpp

375

Seems to me, L is not used in the test, so you can just use {{%.+}} here

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15689–15690

Oops, sorry, wrong comment, thought about Addr_of, but you removed it already.

[cchen]

Fix based on feedback

[ABataev]

Also, would be good to add more tests for the situations that should not be handled, like just (a+b), explicit casts, etc.

clang/include/clang/Basic/DiagnosticSemaKinds.td

10117

I still would change this message to something like expected addressable lvalue in '%0' clause, where %0 must be one of map|to|from to be more specific and correct.

[cchen]

Add back the check of UO_Deref

[cchen]

In D75077#1896708, @cchen wrote:

Add back the check of UO_Deref

Haven't added tests for some prohibited test cases.

[cchen]

Add test

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15500–15502

Shall we do this only if RelevantExpr is nullptr?

15512–15513

Shall we do this only if RelevantExpr is nullptr?

15573–15574

Shall we do this only if RelevantExpr is nullptr?

15681–15685

Shall we do this only if RelevantExpr is nullptr?

15712

Shall we do this only if RelevantExpr is nullptr?

15717

Shall we also here do return RelevantExpr || Visit(LE);?

15718

Shall we also here do return RelevantExpr || Visit(RE);?

[cchen]

clang/lib/Sema/SemaOpenMP.cpp

15500–15502

Yes, I'm assuming that RelevantExpr is nullptr before those assignments, so maybe adding assert is sufficient?

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15500–15502

I think assert will crash for something like x?y:z.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

15500–15502

Though, currently, we just return false in this case. Yes, assert should be enough currently probably.

[cchen]

Add assert to make sure that we never overwrite RelevantExpr

[cchen]

return RelevantExpr || Visit(LE);

[ABataev]

LGTM

[cchen]

@ABataev, can you land it for me when you have time?

Thanks,
Chi Chun

[ABataev]

In D75077#1898678, @cchen wrote:

@ABataev, can you land it for me when you have time?

Thanks,
Chi Chun

Sure

clang/include/clang/Basic/DiagnosticSemaKinds.td

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_flush_order_clause_and_list : Error<
   "'flush' directive with memory order clause '%0' cannot have the list">;
 def note_omp_flush_order_clause_here : Note<
   "memory order clause '%0' is specified here">;
+def err_omp_non_lvalue_in_map_or_motion_clauses: Error<
+  "expected addressable lvalue in '%0' clause"
+  >;
 } // 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

for (; I != CE; ++I) {
       bool IsFinalArraySection =
           isFinalArraySectionExpression(I->getAssociatedExpression());
 
       // Get information on whether the element is a pointer. Have to do a
       // special treatment for array sections given that they are built-in
       // types.
       const auto *OASE =
           dyn_cast<OMPArraySectionExpr>(I->getAssociatedExpression());
+      const auto *UO = dyn_cast<UnaryOperator>(I->getAssociatedExpression());
+      const auto *BO = dyn_cast<BinaryOperator>(I->getAssociatedExpression());
       bool IsPointer =
           (OASE && OMPArraySectionExpr::getBaseOriginalType(OASE)
                        .getCanonicalType()
                        ->isAnyPointerType()) ||
           I->getAssociatedExpression()->getType()->isAnyPointerType();
+      bool IsNonDerefPointer = IsPointer && !UO && !BO;
 
-      if (Next == CE || IsPointer || IsFinalArraySection) {
+      if (Next == CE || IsNonDerefPointer || 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())) &&
+                isa<OMPArraySectionExpr>(Next->getAssociatedExpression()) ||
+                isa<UnaryOperator>(Next->getAssociatedExpression()) ||
+                isa<BinaryOperator>(Next->getAssociatedExpression())) &&
                "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.
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) {
-          const auto *FD = dyn_cast<FieldDecl>(EncounteredME->getMemberDecl());
+          const auto *FD = 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) {

clang/lib/Sema/SemaOpenMP.cpp

class MapBaseChecker final : public StmtVisitor<MapBaseChecker, bool> {
   bool NoDiagnose = false;
   const Expr *RelevantExpr = nullptr;
   bool AllowUnitySizeArraySection = true;
   bool AllowWholeSizeArraySection = true;
   SourceLocation ELoc;
   SourceRange ERange;
 
   void emitErrorMsg() {
-    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)
+    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)
+          << getOpenMPClauseName(CKind) << ERange;
     }
   }
 
 public:
   bool VisitDeclRefExpr(DeclRefExpr *DRE) {
     if (!isa<VarDecl>(DRE->getDecl())) {
       emitErrorMsg();
       return false;
     }
+    assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
     RelevantExpr = DRE;
     // Record the component.
     Components.emplace_back(DRE, DRE->getDecl());
     return true;
   }
 
   bool VisitMemberExpr(MemberExpr *ME) {
     Expr *E = ME;
     Expr *BaseE = ME->getBase()->IgnoreParenCasts();
 
-    if (isa<CXXThisExpr>(BaseE))
+    if (isa<CXXThisExpr>(BaseE)) {
+      assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
       // We found a base expression: this->Val.
       RelevantExpr = ME;
-    else
+    } else {
       E = BaseE;
+    }
 
     if (!isa<FieldDecl>(ME->getMemberDecl())) {
       if (!NoDiagnose) {
         SemaRef.Diag(ELoc, diag::err_omp_expected_access_to_data_field)
           << ME->getSourceRange();
         return false;
       }
       if (RelevantExpr)
bool VisitMemberExpr(MemberExpr *ME) {
     //
     // 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.
     Components.emplace_back(ME, FD);
     return RelevantExpr || Visit(E);
   }
 
   bool VisitArraySubscriptExpr(ArraySubscriptExpr *AE) {
     Expr *E = AE->getBase()->IgnoreParenImpCasts();
 
     if (!E->getType()->isAnyPointerType() && !E->getType()->isArrayType()) {
       if (!NoDiagnose) {
if (const auto *TE = dyn_cast<CXXThisExpr>(E->IgnoreParenCasts())) {
       if (!AE->getIdx()->isValueDependent() &&
           AE->getIdx()->EvaluateAsInt(Result, SemaRef.getASTContext()) &&
           !Result.Val.getInt().isNullValue()) {
         SemaRef.Diag(AE->getIdx()->getExprLoc(),
                      diag::err_omp_invalid_map_this_expr);
         SemaRef.Diag(AE->getIdx()->getExprLoc(),
                      diag::note_omp_invalid_subscript_on_this_ptr_map);
       }
+      assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
       RelevantExpr = TE;
     }
 
     // Record the component - we don't have any declaration associated.
     Components.emplace_back(AE, nullptr);
 
     return RelevantExpr || Visit(E);
   }
if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
           OASE->getLowerBound()->EvaluateAsInt(ResultL,
                                                SemaRef.getASTContext()) &&
           !ResultL.Val.getInt().isNullValue()) {
         SemaRef.Diag(OASE->getLowerBound()->getExprLoc(),
                      diag::err_omp_invalid_map_this_expr);
         SemaRef.Diag(OASE->getLowerBound()->getExprLoc(),
                      diag::note_omp_invalid_lower_bound_on_this_ptr_mapping);
       }
+      assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
       RelevantExpr = TE;
     }
 
     // Record the component - we don't have any declaration associated.
     Components.emplace_back(OASE, nullptr);
     return RelevantExpr || Visit(E);
   }
+  bool VisitUnaryOperator(UnaryOperator *UO) {
+    if (SemaRef.getLangOpts().OpenMP < 50 || !UO->isLValue() ||
+        UO->getOpcode() != UO_Deref) {
+      emitErrorMsg();
+      return false;
+    }
+    if (!RelevantExpr) {
+      // Record the component if haven't found base decl.
+      Components.emplace_back(UO, nullptr);
+    }
+    return RelevantExpr || Visit(UO->getSubExpr()->IgnoreParenImpCasts());
+  }
+  bool VisitBinaryOperator(BinaryOperator *BO) {
+    if (SemaRef.getLangOpts().OpenMP < 50 || !BO->getType()->isPointerType()) {
+      emitErrorMsg();
+      return false;
+    }
+
+    // Pointer arithmetic is the only thing we expect to happen here so after we
+    // make sure the binary operator is a pointer type, the we only thing need
+    // to to is to visit the subtree that has the same type as root (so that we
+    // know the other subtree is just an offset)
+    Expr *LE = BO->getLHS()->IgnoreParenImpCasts();
+    Expr *RE = BO->getRHS()->IgnoreParenImpCasts();
+    Components.emplace_back(BO, nullptr);
+    assert((LE->getType().getTypePtr() == BO->getType().getTypePtr() ||
+            RE->getType().getTypePtr() == BO->getType().getTypePtr()) &&
+           "Either LHS or RHS have base decl inside");
+    if (BO->getType().getTypePtr() == LE->getType().getTypePtr())
+      return RelevantExpr || Visit(LE);
+    return RelevantExpr || Visit(RE);
+  }
+  bool VisitCXXThisExpr(CXXThisExpr *CTE) {
+    assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
+    RelevantExpr = CTE;
+    Components.emplace_back(CTE, nullptr);
+    return true;
+  }
   bool VisitStmt(Stmt *) {
     emitErrorMsg();
     return false;
   }
   const Expr *getFoundBase() const {
     return RelevantExpr;
   }
   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()))
+  if (Checker.Visit(E->IgnoreParens()))
     return Checker.getFoundBase();
   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,
if (VE->isValueDependent() || VE->isTypeDependent() ||
       // We can only analyze this information once the missing information is
       // resolved.
       MVLI.ProcessedVarList.push_back(RE);
       continue;
     }
 
     Expr *SimpleExpr = RE->IgnoreParenCasts();
 
-    if (!RE->IgnoreParenImpCasts()->isLValue()) {
-      SemaRef.Diag(ELoc,
-                   diag::err_omp_expected_named_var_member_or_array_expression)
+    if (!RE->isLValue()) {
+      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)
+            << getOpenMPClauseName(CKind) << 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.

clang/test/OpenMP/target_map_messages.cpp

-// RUN: %clang_cc1 -verify=expected,warn -fopenmp -ferror-limit 200 %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
 // RUN: %clang_cc1 -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-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=40 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
 // RUN: %clang_cc1 -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 addressable lvalue in 'map' clause}}
     {}
-    #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 addressable lvalue in 'map' clause}}
   {}
   #pragma omp target map(r.C, r.D)
   {}
   #pragma omp target map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
   {}
   #pragma omp target map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
   {}
   #pragma omp target map(r.C, r.S)  // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
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 addressable lvalue in 'map' clause}}
   foo();
 #pragma omp target map(S2::S2s)
   foo();
 #pragma omp target map(S2::S2sc)
   foo();
 #pragma omp target map(x)
   foo();
 #pragma omp target map(to: x)
   foo();
 #pragma omp target map(to: to)
   foo();
 #pragma omp target map(to)
   foo();
 #pragma omp target map(to, x)
   foo();
 #pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
-#pragma omp target data map(tofrom: argc > 0 ? x : y) // 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 addressable lvalue in 'map' clause}}
 #pragma omp target data map(argc)
 #pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
 #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} warn-warning 2 {{Type 'const S2' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning 2 {{Type 'const S3' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(ba) // warn-warning 2 {{Type 'const S2 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(ca) // warn-warning 2 {{Type 'const S3 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)
 #pragma omp target data map(S2::S2sc)
#pragma omp target 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 addressable lvalue in 'map' clause}}
 #pragma omp target data map(argc)
 #pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
 #pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} warn-warning {{Type 'const S2' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'const S3' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(argv[1])
 #pragma omp target data map(ba) // warn-warning {{Type 'const S2 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(ca) // warn-warning {{Type 'const S3 [5]' is not trivially copyable and not guaranteed to be mapped correctly}}
 #pragma omp target data map(da)
 #pragma omp target data map(S2::S2s)

 #pragma omp target map(s.p->p->p->a)
 // expected-error@+1 {{variable already marked as mapped in current construct}}
   { s.a++; }
 #pragma omp target map(s.s.s.b[:2])
   { s.s.s.b[0]++; }
 #pragma omp target map(s8[0:1], s9) // warn-warning {{Type 'class S8' is not trivially copyable and not guaranteed to be mapped correctly}} warn-warning {{Type 'class S9' is not trivially copyable and not guaranteed to be mapped correctly}}
   {}
 
+  int **BB, *offset, *a;
+
+#pragma omp target map(**(BB+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target map(**(BB+y)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target map(*(a+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target map(**(*offset+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target map(**(y+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target map(*(*offset+a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target map(**(*offset+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target map(**(*(*(&offset))+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target map(*(a+(a))) // expected-error {{invalid operands to binary expression ('int *' and 'int *')}}
+  {}
+#pragma omp target map(*(1+*a+*a)) // expected-error {{indirection requires pointer operand ('int' invalid)}}
+  {}
+
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 #endif

clang/test/OpenMP/target_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,le50 -fopenmp-version=50 -fopenmp -ferror-limit 200 %s -Wno-openmp-mapping -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
 #ifdef CCODE
 void foo(int arg) {
   const int n = 0;
 
   double marr[10][10][10];
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 addressable lvalue in 'map' clause}} 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 addressable lvalue in 'map' clause}} 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 addressable lvalue in 'map' clause}} 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 addressable lvalue in 'map' clause}} 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 addressable lvalue in 'map' clause}} 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)
#pragma omp target teams private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}}  le45-note {{defined as private}}
   {}
 
 #pragma omp target teams firstprivate(j) map(j) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} le45-note {{defined as firstprivate}}
   {}
 
 #pragma omp target teams map(m)
   {}
 
+  int **BB, *offset, *a;
+
+#pragma omp target teams map(**(BB+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target teams map(**(BB+y)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target teams map(*(a+*offset)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target teams map(**(*offset+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target teams map(**(y+BB)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target teams map(*(*offset+a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target teams map(**(*offset+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target teams map(**(*(*(&offset))+BB+*a)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+#pragma omp target teams map(*(a+(a))) // expected-error {{invalid operands to binary expression ('int *' and 'int *')}}
+  {}
+#pragma omp target teams map(*(1+*a+*a)) // expected-error {{indirection requires pointer operand ('int' invalid)}}
+  {}
+
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 #endif

clang/test/OpenMP/target_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* [[B_VAL:%.+]], i32** [[BPC0]]
+  // CK5-DAG: store i32* [[B_VAL_2:%.+]], i32** [[PC0]]
+  // CK5-DAG: [[B_VAL]] = load i32*, i32** [[B_ADDR:%.+]]
+  // CK5-DAG: [[B_VAL_2]] = 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-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i{{32|64}} [[IDX_EXT:%.+]]
+  // CK6-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i{{32|64}} [[L_VAL:%.+]]
+  // CK6-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
+  // CK6-DAG: [[L_VAL]] = load i32, i32* [[L_ADDR:%.+]]
+  // CK6-DAG: store i32 {{.+}}, 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* [[B_VAL:%.+]], i32** [[BPC0]]
+  // CK7-DAG: store i32* [[ARRAY_IDX:%.+]], i32** [[PC0]]
+  // CK7-DAG: [[B_VAL]] = 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* [[B_VAL_2:%.+]], i32 [[L_VAL:%.+]]
+  // CK7-32-DAG: [[B_VAL_2]] = load i32*, i32** [[B_ADDR]]
+  // CK7-32-DAG: [[L_VAL]] = load i32, i32* [[L_ADDR:%.+]]
+  // CK7-32-DAG: [[IDX_PROM]] = load i32, i32* [[L_ADDR]]
+  // CK7-64-DAG: [[IDX_EXT:%.+]] = sext i32 [[L_VAL:%.+]] to i64
+  // CK7-64-DAG: [[L_VAL:%.+]] = load i32, i32* [[L_ADDR:%.+]]
+  // CK7-64-DAG: [[IDX_PROM]] = sext i32 [[L_VAL_2:%.+]] to i64
+  // CK7-64-DAG: [[L_VAL_2]] = 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* [[P_VAL:%.+]], double** [[PC0]]
+  // CK9-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE0]]
+  // CK9-DAG: [[P]] = getelementptr inbounds [[STRUCT_S:%.+]], [[STRUCT_S]]* [[S_VAL:%.+]], i32 0, i32 0
+  // CK9-DAG: [[S_VAL]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S_ADDR:%.+]]
+  // CK9-DAG: [[P_VAL]] = load double*, double** [[P_1:%.+]],
+  // CK9-DAG: [[P_1]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[S_VAL_2:%.+]], i32 0, i32 0
+  // CK9-DAG: [[S_VAL_2]] = 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* [[S_P:%.+]], i{{.+}} [[IDX_EXT:%.+]]
+  // CK10-32-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[S_P:%.+]], i{{.+}} [[L_VAL:%.+]]
+  // CK10-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
+  // CK10-DAG: [[S_P]] = 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* [[S_P:%.+]], i{{.+}} [[IDX_EXT:%.+]]
+  // CK11-32-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[S_P:%.+]], i{{.+}} [[L_VAL:%.+]]
+  // CK11-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
+  // CK11-DAG: [[S_P]] = load double*, double** [[P_1:%.+]],
+  // CK11-DAG: [[P_1]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[S_ADDR:%.+]], i32 0, i32 0
+  #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
+///==========================================================================///
+// RUN: %clang_cc1 -DCK13 -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 CK13 --check-prefix CK13-64
+// RUN: %clang_cc1 -DCK13 -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 CK13 --check-prefix CK13-64
+// RUN: %clang_cc1 -DCK13 -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 CK13 --check-prefix CK13-32
+// RUN: %clang_cc1 -DCK13 -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 CK13 --check-prefix CK13-32
+
+// RUN: %clang_cc1 -DCK13 -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 -DCK13 -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 -DCK13 -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 -DCK13 -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 CK13
+
+// CK13: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK13: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+// CK13-LABEL: lvalue
+void lvalue(int **BB, int a, int b) {
+
+  // CK13-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]]{{.+}})
+  // CK13-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK13-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK13-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32***
+  // CK13-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK13-DAG: store i32** [[ZERO:%.+]], i32*** [[BPC0]]
+  // CK13-DAG: store i32* [[ADD_PTR_2:%.+]], i32** [[PC0]]
+  // CK13-DAG: [[ZERO]] = load i32**, i32*** [[BB_ADDR:%.+]]
+  // CK13-64-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32, i32* [[RESULT:%.+]], i64 [[IDX_EXT_1:%.+]]
+  // CK13-32-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32, i32* [[RESULT:%.+]], i32 [[B_ADDR:%.+]]
+  // CK13-64-DAG: [[IDX_EXT_1]] = sext i32 [[B_ADDR:%.+]]
+  // CK13-DAG: [[RESULT]] = load i32*, i32** [[ADD_PTR:%.+]],
+  // CK13-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i64 [[IDX_EXT:%.+]]
+  // CK13-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i32 [[A_ADDR:%.+]]
+  // CK13-64-DAG: [[IDX_EXT]] = sext i32 [[TWO:%.+]] to i64
+  // CK13-DAG: [[B_VAL]] = load i32**, i32*** [[BB_ADDR]]
+  #pragma omp target update to(*(*(BB+a)+b))
+  *(*(BB+a)+b) = 1;
+  #pragma omp target update from(*(*(BB+a)+b))
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK14 -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 CK14 --check-prefix CK14-64
+// RUN: %clang_cc1 -DCK14 -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 CK14 --check-prefix CK14-64
+// RUN: %clang_cc1 -DCK14 -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 CK14 --check-prefix CK14-32
+// RUN: %clang_cc1 -DCK14 -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 CK14 --check-prefix CK14-32
+
+// RUN: %clang_cc1 -DCK14 -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 -DCK14 -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 -DCK14 -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 -DCK14 -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 CK14
+
+// CK14: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710673]
+
+struct SSA {
+  double *p;
+  double *≺
+  SSA(double *&pr) : pr(pr) {}
+};
+
+struct SSB {
+  double *d;
+  SSA *p;
+  SSA *≺
+  SSB(SSA *&pr) : pr(pr) {}
+
+  // CK14-LABEL: define {{.+}}foo
+  void foo() {
+
+    // CK14-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+    // CK14-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+    // CK14-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+    // CK14-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+
+    // CK14-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+    // CK14-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+    // CK14-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+    // CK14-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to double***
+    // CK14-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double**
+    // CK14-DAG: store double** [[D_VAL:%.+]], double*** [[BPC1]]
+    // CK14-DAG: store double* [[ADD_PTR:%.+]], double** [[PC1]]
+    // CK14-DAG: store i64 8, i64* [[SIZE1]]
+    // CK14-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[ZERO:%.+]], i{{.+}} 1
+    // CK14-DAG: [[ZERO]] = load double*, double** [[D_VAL_2:%.+]]
+    // CK14-DAG: [[D_VAL]] = getelementptr inbounds [[SSB:%.+]], [[SSB:%.+]]* [[THIS:%.+]], i32 0, i32 0
+    // CK14-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+    // CK14-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+    // CK14-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0
+    // CK14-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[SSB]]**
+    // CK14-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double***
+    // CK14-DAG: store [[SSB]]* [[THIS]], [[SSB]]** [[BPC0]],
+    // CK14-DAG: store double** [[D_VAL]], double*** [[PC0]]
+    // CK14-DAG: store i{{.+}} [[COMPUTE_LEN:%.+]], i{{.+}}* [[SIZE0]]
+
+    #pragma omp target update to(*(this->d+1))
+    *(this->d+1) = 1;
+    #pragma omp target update from(*(this->d+1))
+  }
+};
+
+void lvalue_member(SSA *sap) {
+  double *d;
+  SSA sa(d);
+  SSB sb(sap);
+  sb.foo();
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK15 -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 CK15 --check-prefix CK15-64
+// RUN: %clang_cc1 -DCK15 -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 CK15 --check-prefix CK15-64
+// RUN: %clang_cc1 -DCK15 -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 CK15 --check-prefix CK15-32
+// RUN: %clang_cc1 -DCK15 -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 CK15 --check-prefix CK15-32
+
+// RUN: %clang_cc1 -DCK15 -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 -DCK15 -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 -DCK15 -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 -DCK15 -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 CK15
+
+// CK15: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673]
+
+struct SSA {
+  double *p;
+  double *≺
+  SSA(double *&pr) : pr(pr) {}
+};
+
+//CK-15-LABEL: lvalue_member
+void lvalue_member(SSA *sap) {
+
+  // CK15-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK15-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK15-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK15-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+
+  // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK15-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK15-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to double***
+  // CK15-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double**
+  // CK15-DAG: store double** [[P_VAL:%.+]], double*** [[BPC1]]
+  // CK15-DAG: store double* [[ADD_PTR:%.+]], double** [[PC1]]
+  // CK15-DAG: store i64 8, i64* [[SIZE1]]
+  // CK15-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} 3
+  // CK15-DAG: [[THREE]] = load double*, double** [[P_VAL_1:%.+]]
+  // CK15-DAG: [[P_VAL]] = getelementptr inbounds [[SSA:%.+]], [[SSA:%.+]]* [[THIS:%.+]], i32 0, i32 0
+  // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK15-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0
+  // CK15-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[SSA]]**
+  // CK15-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double***
+  // CK15-DAG: store [[SSA]]* [[SAP_VAL:%.+]], [[SSA]]** [[BPC0]],
+  // CK15-DAG: store double** [[P_VAL]], double*** [[PC0]]
+  // CK15-DAG: store i{{.+}} [[COMPUTE_SIZE:%.+]], i{{.+}}* [[SIZE0]]
+  // CK15-DAG: [[SAP_VAL]] = load [[SSA]]*, [[SSA]]** [[SAP_ADDR:%.+]],
+  // CK15-DAG: [[COMPUTE_SIZE]] = sdiv exact i64 [[NINE:%.+]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64)
+  // CK15-DAG: [[NINE]] = sub i{{.+}} [[SEVEN:%.+]], [[EIGHT:%.+]]
+  // CK15-DAG: [[SEVEN]] = ptrtoint i8* [[SIX:%.+]] to i64
+  // CK15-DAG: [[EIGHT]] = ptrtoint i8* [[FIVE:%.+]] to i64
+  // CK15-DAG: [[SIX]] = bitcast double** {{.+}} to i8*
+  // CK15-DAG: [[FIVE]] = bitcast double** {{.+}} to i8*
+  #pragma omp target update to(*(3+sap->p))
+  *(3+sap->p) = 1;
+  #pragma omp target update from(*(3+sap->p))
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK16 -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 CK16 --check-prefix CK16-64
+// RUN: %clang_cc1 -DCK16 -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 CK16 --check-prefix CK16-64
+// RUN: %clang_cc1 -DCK16 -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 CK16 --check-prefix CK16-32
+// RUN: %clang_cc1 -DCK16 -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 CK16 --check-prefix CK16-32
+
+// RUN: %clang_cc1 -DCK16 -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 -DCK16 -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 -DCK16 -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 -DCK16 -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 CK16
+
+// CK16: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK16: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+//CK16-LABEL: lvalue_find_base
+void lvalue_find_base(float *f, int *i) {
+
+  // CK16-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]]{{.+}})
+  // CK16-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK16-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK16-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK16-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK16-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**
+  // CK16-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK16-DAG: store float* [[ZERO:%.+]], float** [[BPC0]]
+  // CK16-DAG: store float* [[ADD_PTR:%.+]], float** [[PC0]]
+  // CK16-32-DAG: [[ADD_PTR]] = getelementptr inbounds float, float* [[THREE:%.+]], i32 [[I:%.+]]
+  // CK16-64-DAG: [[ADD_PTR]] = getelementptr inbounds float, float* [[THREE:%.+]], i64 [[IDX_EXT:%.+]]
+  // CK16-DAG: [[THREE]] = load float*, float** [[F_ADDR:%.+]],
+  // CK16-64-DAG: [[IDX_EXT]] = sext i32 [[I:%.+]] to i64
+  // CK16-DAG: [[ZERO]] = load float*, float** [[F_ADDR:%.+]]
+
+  #pragma omp target update to(*(*i+f))
+  *(*i+f) = 1.0;
+  #pragma omp target update from(*(*i+f))
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK17 -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 CK17 --check-prefix CK17-64
+// RUN: %clang_cc1 -DCK17 -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 CK17 --check-prefix CK17-64
+// RUN: %clang_cc1 -DCK17 -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 CK17 --check-prefix CK17-32
+// RUN: %clang_cc1 -DCK17 -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 CK17 --check-prefix CK17-32
+
+// RUN: %clang_cc1 -DCK17 -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 -DCK17 -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 -DCK17 -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 -DCK17 -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 CK17
+
+// CK17: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK17: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+struct SSA {
+  int i;
+  SSA *sa;
+};
+
+//CK17-LABEL: lvalue_find_base
+void lvalue_find_base(float **f, SSA *sa) {
+
+  // CK17-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]]{{.+}})
+  // CK17-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK17-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK17-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK17-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK17-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float***
+  // CK17-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float**
+  // CK17-DAG: store float** [[ZERO:%.+]], float*** [[BPC0]],
+  // CK17-DAG: store float* [[ADD_PTR_4:%.+]], float** [[PC0]],
+  // CK17-DAG: [[ZERO]] = load float**, float*** [[F_ADDR:%.+]],
+  // CK17-64-DAG: [[ADD_PTR_4]] = getelementptr inbounds float, float* [[SEVEN:%.+]], i64 [[IDX_EXT_3:%.+]]
+  // CK17-64-DAG: [[IDX_EXT_3]] = sext i32 [[I_VAL:%.+]] to i64
+  // CK17-32-DAG: [[ADD_PTR_4]] = getelementptr inbounds float, float* [[SEVEN:%.+]], i32 [[I_VAL:%.+]]
+  // CK17-DAG: [[I_VAL]] = load i32, i32* [[I:%.+]],
+  // CK17-DAG: [[SEVEN]] = load float*, float** [[ADD_PTR:%.+]],
+  // CK17-64-DAG: [[ADD_PTR]] = getelementptr inbounds float*, float** [[F:%.+]], i64 [[IDX_EXT:%.+]]
+  // CK17-32-DAG: [[ADD_PTR]] = getelementptr inbounds float*, float** [[F:%.+]], i32 [[ADD:%.+]]
+  // CK17-64-DAG: [[IDX_EXT]] = sext i32 [[ADD:%.+]] to i64
+  // CK17-DAG: [[ADD]] = add nsw i32 1, [[FIVE:%.+]]
+  // CK17-DAG: [[FIVE]] = load i32, i32* [[I_2:%.+]],
+  // CK17-DAG: [[I_2]] = getelementptr inbounds [[SSA:%.+]], [[SSA]]* [[FOUR:%.+]], i32 0, i32 0
+  // CK17-DAG: [[FOUR]] = load [[SSA]]*, [[SSA]]** [[SSA_ADDR:%.+]],
+
+  #pragma omp target update to(*(sa->sa->i+*(1+sa->i+f)))
+  *(sa->sa->i+*(1+sa->i+f)) = 1;
+  #pragma omp target update from(*(sa->sa->i+*(1+sa->i+f)))
+}
+
+#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->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'}}
+#pragma omp target update from(*(this+this)) // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} expected-error {{invalid operands to binary expression ('S8 *' and 'S8 *')}}
+  }
+};
 
 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 addressable lvalue in 'from' clause}} 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 addressable lvalue in 'from' clause}} 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)
int main(int argc, char **argv) {
   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 addressable lvalue in 'from' clause}} 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(**(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))) // 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+(m))) // expected-error {{invalid operands to binary expression ('int *' and 'int *')}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(*(1+y+y)) // expected-error {{indirection requires pointer operand ('int' invalid)}} expected-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)
   }
 
   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->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'}}
+#pragma omp target update to(*(this+this)) // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} expected-error {{invalid operands to binary expression ('S8 *' and 'S8 *')}}
+    {}
+  }
+};
 
 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 addressable lvalue in 'to' clause}} 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 addressable lvalue in 'to' clause}} 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)
int main(int argc, char **argv) {
   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 addressable lvalue in 'to' clause}} 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(**(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))) // 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+(m))) // expected-error {{invalid operands to binary expression ('int *' and 'int *')}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(*(1+y+y)) // expected-error {{indirection requires pointer operand ('int' invalid)}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {}
   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}}
 }