diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
--- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp
+++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp
@@ -7406,6 +7406,7 @@
     //   S1 s;
     //   double *p;
     //   struct S2 *ps;
+    //   struct S2 **ps_ptr;
     // }
     // S2 s;
     // S2 *ps;
@@ -7542,6 +7543,48 @@
     // &(ps->p), &(ps->p[0]), 33*sizeof(double), MEMBER_OF(4) | PTR_AND_OBJ | TO
     // (*) the struct this entry pertains to is the 4th element in the list
     //     of arguments, hence MEMBER_OF(4)
+    //
+    // map(*p)
+    // &p, &(p[0]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map(*(p+3))
+    // &p, &(p[3]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map((p+3)+7)
+    // &p, &(p[10]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map(*(ps->p))
+    // ps, &(ps->p[0]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map(*(ps->p+3))
+    // ps, &(ps->p[3]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map((ps->p+3)[1])
+    // ps, &(ps->p[4]), 1*sizeof(float), TARGET_PARAM | TO | FROM
+    //
+    // map(*(ps->ps->p))
+    // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM
+    // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1)
+    // &(ps->ps), &(ps->ps->p[0]), sizeof(double), MEMBER_OF(1) | PTR_AND_OBJ |
+    //  FROM
+    //
+    // map((ps->ps->p)[3])
+    // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM
+    // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1)
+    // &(ps->ps), &(ps->ps->p[3]), sizeof(double), MEMBER_OF(1) | PTR_AND_OBJ |
+    //  FROM
+    //
+    // map(*(ps->ps->p+3))
+    // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM
+    // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1)
+    // &(ps->ps), &(ps->ps->p[3]), sizeof(double), MEMBER_OF(1) | PTR_AND_OBJ |
+    //  FROM
+    //
+    // map(*(ps->ps->ps_ptr))
+    // ps, &(ps->ps), sizeof(S2*), TARGET_PARAM
+    // ps, &(ps->ps), sizeof(S2*), MEMBER_OF(1)
+    // &(ps->ps_ptr), &(ps->ps->ps_ptr[0]), sizeof(S2*), MEMBER_OF(1) |
+    //   PTR_AND_OBJ | TO | FROM
 
     // Track if the map information being generated is the first for a capture.
     bool IsCaptureFirstInfo = IsFirstComponentList;
@@ -7665,7 +7708,9 @@
         assert((Next == CE ||
                 isa<MemberExpr>(Next->getAssociatedExpression()) ||
                 isa<ArraySubscriptExpr>(Next->getAssociatedExpression()) ||
-                isa<OMPArraySectionExpr>(Next->getAssociatedExpression())) &&
+                isa<OMPArraySectionExpr>(Next->getAssociatedExpression()) ||
+                isa<UnaryOperator>(Next->getAssociatedExpression()) ||
+                isa<ConditionalOperator>(Next->getAssociatedExpression())) &&
                "Unexpected expression");
 
         Address LB = CGF.EmitOMPSharedLValue(I->getAssociatedExpression())
@@ -7781,18 +7826,21 @@
         // mapped member. If the parent is "*this", then the value declaration
         // is nullptr.
         if (EncounteredME) {
-          const auto *FD = dyn_cast<FieldDecl>(EncounteredME->getMemberDecl());
-          unsigned FieldIndex = FD->getFieldIndex();
-
-          // Update info about the lowest and highest elements for this struct
-          if (!PartialStruct.Base.isValid()) {
-            PartialStruct.LowestElem = {FieldIndex, LB};
-            PartialStruct.HighestElem = {FieldIndex, LB};
-            PartialStruct.Base = BP;
-          } else if (FieldIndex < PartialStruct.LowestElem.first) {
-            PartialStruct.LowestElem = {FieldIndex, LB};
-          } else if (FieldIndex > PartialStruct.HighestElem.first) {
-            PartialStruct.HighestElem = {FieldIndex, LB};
+          // MemberExpr can also be FunctionDecl after we allow all lvalue
+          if (const auto *FD =
+                  dyn_cast<FieldDecl>(EncounteredME->getMemberDecl())) {
+            unsigned FieldIndex = FD->getFieldIndex();
+
+            // Update info about the lowest and highest elements for this struct
+            if (!PartialStruct.Base.isValid()) {
+              PartialStruct.LowestElem = {FieldIndex, LB};
+              PartialStruct.HighestElem = {FieldIndex, LB};
+              PartialStruct.Base = BP;
+            } else if (FieldIndex < PartialStruct.LowestElem.first) {
+              PartialStruct.LowestElem = {FieldIndex, LB};
+            } else if (FieldIndex > PartialStruct.HighestElem.first) {
+              PartialStruct.HighestElem = {FieldIndex, LB};
+            }
           }
         }
 
diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp
--- a/clang/lib/Sema/SemaOpenMP.cpp
+++ b/clang/lib/Sema/SemaOpenMP.cpp
@@ -15197,6 +15197,94 @@
   return ConstLength.getSExtValue() != 1;
 }
 
+namespace {
+class LocatorChecker final : public StmtVisitor<LocatorChecker, bool> {
+  OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents;
+  DeclRefExpr *DRE;
+  // Base decl must be pointer type if we seen unary operator, array
+  // subscriptor, or array section expr, and if we found more than one
+  // DeclRefExpr, we set the base decl as the one being pointer type
+  size_t DerefCnt;
+
+public:
+  bool VisitDeclRefExpr(DeclRefExpr *E) {
+    DRE = E;
+    const auto *VD = dyn_cast<VarDecl>(E->getDecl());
+    const auto *FD = dyn_cast<FieldDecl>(E->getDecl());
+    const auto *FunD = dyn_cast<FunctionDecl>(E->getDecl());
+    if (VD || FD || FunD) {
+      const clang::Type *PT = E->getType().getTypePtr();
+      const std::string TypeStr = PT->getCanonicalTypeInternal().getAsString();
+      size_t PtrCnt = std::count(TypeStr.begin(), TypeStr.end(), '*');
+      if (DerefCnt == PtrCnt) {
+        CurComponents.emplace_back(E, E->getDecl());
+      }
+      return true;
+    }
+    return false;
+  }
+  bool VisitMemberExpr(MemberExpr *ME) {
+    if (auto *FD = dyn_cast<FieldDecl>(ME->getMemberDecl())) {
+      CurComponents.emplace_back(ME, FD->getCanonicalDecl());
+      Visit(ME->getBase()->IgnoreParenImpCasts());
+      return true;
+    }
+    return false;
+  }
+  bool VisitUnaryOperator(UnaryOperator *UO) {
+    DerefCnt++;
+    CurComponents.emplace_back(UO, nullptr);
+    Expr *E = UO->getSubExpr()->IgnoreParenImpCasts();
+    if (E && Visit(E)) {
+      DerefCnt--;
+      return true;
+    }
+    return false;
+  }
+  // checkMapClauseExpressionBase has already added the CurComponents, and we
+  // didn't clean it up, so don't add it here
+  bool VisitArraySubscriptExpr(ArraySubscriptExpr *AE) {
+    DerefCnt++;
+    Expr *E = AE->getBase()->IgnoreParenImpCasts();
+    if (E && Visit(E)) {
+      DerefCnt--;
+      return true;
+    }
+    return false;
+  }
+  // checkMapClauseExpressionBase has already added the CurComponents, and we
+  // didn't clean it up, so don't add it here
+  bool VisitOMPArraySectionExpr(OMPArraySectionExpr *OASE) {
+    DerefCnt++;
+    Expr *E = OASE->getBase()->IgnoreParenImpCasts();
+    if (E && Visit(E)) {
+      DerefCnt--;
+      return true;
+    }
+    return false;
+  }
+  bool VisitBinaryOperator(BinaryOperator *BO) {
+    Expr *LE = BO->getLHS()->IgnoreParenImpCasts();
+    Expr *RE = BO->getRHS()->IgnoreParenImpCasts();
+    if (LE && Visit(LE) && RE && Visit(RE)) {
+      return true;
+    }
+    return false;
+  }
+  bool VisitStmt(Stmt *S) {
+    for (Stmt *Child : S->children()) {
+      if (Child && Visit(Child))
+        return true;
+    }
+    return false;
+  }
+  DeclRefExpr *GetDRE() const { return DRE; }
+  explicit LocatorChecker(
+      OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents)
+      : CurComponents(CurComponents), DerefCnt(0) {}
+};
+} // 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
@@ -15249,11 +15337,14 @@
   bool AllowUnitySizeArraySection = true;
   bool AllowWholeSizeArraySection = true;
 
+  Expr *OrigExpr = E;
+
   while (!RelevantExpr) {
     E = E->IgnoreParenImpCasts();
 
     if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
-      if (!isa<VarDecl>(CurE->getDecl()))
+      if (!(isa<VarDecl>(CurE->getDecl()) ||
+            isa<FunctionDecl>(CurE->getDecl())))
         return nullptr;
 
       RelevantExpr = CurE;
@@ -15274,7 +15365,8 @@
       else
         E = BaseE;
 
-      if (!isa<FieldDecl>(CurE->getMemberDecl())) {
+      if (!isa<FieldDecl>(CurE->getMemberDecl()) &&
+          !isa<CXXMethodDecl>(CurE->getMemberDecl())) {
         if (!NoDiagnose) {
           SemaRef.Diag(ELoc, diag::err_omp_expected_access_to_data_field)
               << CurE->getSourceRange();
@@ -15285,12 +15377,13 @@
         continue;
       }
 
-      auto *FD = cast<FieldDecl>(CurE->getMemberDecl());
+      auto *FD = dyn_cast<FieldDecl>(CurE->getMemberDecl());
+      auto *FunD = dyn_cast<FunctionDecl>(CurE->getMemberDecl());
 
       // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.3]
       //  A bit-field cannot appear in a map clause.
       //
-      if (FD->isBitField()) {
+      if (FD && FD->isBitField()) {
         if (!NoDiagnose) {
           SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_clause)
               << CurE->getSourceRange() << getOpenMPClauseName(CKind);
@@ -15330,7 +15423,10 @@
       AllowWholeSizeArraySection = false;
 
       // Record the component.
-      CurComponents.emplace_back(CurE, FD);
+      if (FD)
+        CurComponents.emplace_back(CurE, FD);
+      else
+        CurComponents.emplace_back(CurE, FunD);
     } else if (auto *CurE = dyn_cast<ArraySubscriptExpr>(E)) {
       E = CurE->getBase()->IgnoreParenImpCasts();
 
@@ -15436,6 +15532,12 @@
       // Record the component - we don't have any declaration associated.
       CurComponents.emplace_back(CurE, nullptr);
     } else {
+      if (OrigExpr->IgnoreParenImpCasts()->isLValue()) {
+        LocatorChecker Checker(CurComponents);
+        if (Checker.Visit(OrigExpr)) {
+          return Checker.GetDRE();
+        }
+      }
       if (!NoDiagnose) {
         // If nothing else worked, this is not a valid map clause expression.
         SemaRef.Diag(
@@ -15898,8 +16000,11 @@
     }
 
     Expr *SimpleExpr = RE->IgnoreParenCasts();
+    clang::Expr::isModifiableLvalueResult MLValue =
+        RE->isModifiableLvalue(SemaRef.getASTContext());
 
-    if (!RE->IgnoreParenImpCasts()->isLValue()) {
+    if (!RE->IgnoreParenImpCasts()->isLValue() &&
+        (MLValue != clang::Expr::isModifiableLvalueResult::MLV_LValueCast)) {
       SemaRef.Diag(ELoc,
                    diag::err_omp_expected_named_var_member_or_array_expression)
           << RE->getSourceRange();
@@ -15949,8 +16054,10 @@
 
     auto *VD = dyn_cast<VarDecl>(CurDeclaration);
     const auto *FD = dyn_cast<FieldDecl>(CurDeclaration);
+    const auto *FuncD = dyn_cast<FunctionDecl>(CurDeclaration);
 
-    assert((VD || FD) && "Only variables or fields are expected here!");
+    assert((VD || FD || FuncD) &&
+           "Only variables, fields, or function calls are expected here!");
     (void)FD;
 
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.10]
diff --git a/clang/test/OpenMP/target_map_messages.cpp b/clang/test/OpenMP/target_map_messages.cpp
--- a/clang/test/OpenMP/target_map_messages.cpp
+++ b/clang/test/OpenMP/target_map_messages.cpp
@@ -59,7 +59,7 @@
     {}
     #pragma omp target map(arg,a*2) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
     {}
-    #pragma omp target map(arg,(c+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    #pragma omp target map(arg,(c+1)[2])
     {}
     #pragma omp target map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}}
     {}
@@ -305,7 +305,7 @@
   {}
   #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])
   {}
   #pragma omp target map(r.C, r.D)
   {}
@@ -339,7 +339,7 @@
   {}
   #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)
   {}
   #pragma omp target map(u.B)  // expected-error {{mapping of union members is not allowed}}
   {}
@@ -496,7 +496,7 @@
 #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)
 #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}}
@@ -592,7 +592,7 @@
 #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])
 #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}}
diff --git a/clang/test/OpenMP/target_parallel_for_map_messages.cpp b/clang/test/OpenMP/target_parallel_for_map_messages.cpp
--- a/clang/test/OpenMP/target_parallel_for_map_messages.cpp
+++ b/clang/test/OpenMP/target_parallel_for_map_messages.cpp
@@ -121,7 +121,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for 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 parallel for map(tofrom: argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(argc)
   for (i = 0; i < argc; ++i) foo();
@@ -229,7 +229,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target parallel for map(tofrom: argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for map(argc)
   for (i = 0; i < argc; ++i) foo();
diff --git a/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp b/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
--- a/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
+++ b/clang/test/OpenMP/target_parallel_for_simd_map_messages.cpp
@@ -121,7 +121,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd 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 parallel for simd map(tofrom: argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(argc)
   for (i = 0; i < argc; ++i) foo();
@@ -229,7 +229,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target parallel for simd map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target parallel for simd map(tofrom: argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target parallel for simd map(argc)
   for (i = 0; i < argc; ++i) foo();
diff --git a/clang/test/OpenMP/target_parallel_map_messages.cpp b/clang/test/OpenMP/target_parallel_map_messages.cpp
--- a/clang/test/OpenMP/target_parallel_map_messages.cpp
+++ b/clang/test/OpenMP/target_parallel_map_messages.cpp
@@ -121,7 +121,7 @@
   foo();
 #pragma omp target parallel map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
-#pragma omp target parallel 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 parallel map(tofrom: argc > 0 ? x : y)
   foo();
 #pragma omp target parallel map(argc)
   foo();
@@ -228,7 +228,7 @@
   foo();
 #pragma omp target parallel map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
-#pragma omp target parallel map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target parallel map(tofrom: argc > 0 ? argv[1] : argv[2])
   foo();
 #pragma omp target parallel map(argc)
   foo();
diff --git a/clang/test/OpenMP/target_simd_map_messages.cpp b/clang/test/OpenMP/target_simd_map_messages.cpp
--- a/clang/test/OpenMP/target_simd_map_messages.cpp
+++ b/clang/test/OpenMP/target_simd_map_messages.cpp
@@ -115,7 +115,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd 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 simd map(tofrom: argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(argc)
   for (i = 0; i < argc; ++i) foo();
@@ -217,7 +217,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target simd map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target simd map(tofrom: argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target simd map(argc)
   for (i = 0; i < argc; ++i) foo();
diff --git a/clang/test/OpenMP/target_teams_distribute_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_distribute_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_map_messages.cpp
@@ -121,7 +121,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute 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 teams distribute map(tofrom: argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(argc)
   for (i = 0; i < argc; ++i) foo();
@@ -229,7 +229,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams distribute map(tofrom: argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute map(argc)
   for (i = 0; i < argc; ++i) foo();
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_map_messages.cpp
@@ -121,7 +121,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for 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 teams distribute parallel for map(tofrom: argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(argc)
   for (i = 0; i < argc; ++i) foo();
@@ -229,7 +229,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams distribute parallel for map(tofrom: argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for map(argc)
   for (i = 0; i < argc; ++i) foo();
diff --git a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_parallel_for_simd_map_messages.cpp
@@ -121,7 +121,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd 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 teams distribute parallel for simd map(tofrom: argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(argc)
   for (i = 0; i < argc; ++i) foo();
@@ -229,7 +229,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute parallel for simd map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams distribute parallel for simd map(tofrom: argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute parallel for simd map(argc)
   for (i = 0; i < argc; ++i) foo();
diff --git a/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp b/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_distribute_simd_map_messages.cpp
@@ -121,7 +121,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd 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 teams distribute simd map(tofrom: argc > 0 ? x : y)
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(argc)
   for (i = 0; i < argc; ++i) foo();
@@ -229,7 +229,7 @@
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
   for (i = 0; i < argc; ++i) foo();
-#pragma omp target teams distribute simd map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target teams distribute simd map(tofrom: argc > 0 ? argv[1] : argv[2])
   for (i = 0; i < argc; ++i) foo();
 #pragma omp target teams distribute simd map(argc)
   for (i = 0; i < argc; ++i) foo();
diff --git a/clang/test/OpenMP/target_teams_map_messages.cpp b/clang/test/OpenMP/target_teams_map_messages.cpp
--- a/clang/test/OpenMP/target_teams_map_messages.cpp
+++ b/clang/test/OpenMP/target_teams_map_messages.cpp
@@ -46,7 +46,7 @@
     {}
     #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,(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])
     {}
     #pragma omp target teams map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}}
     {}
@@ -252,7 +252,7 @@
   {}
   #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])
   {}
   #pragma omp target teams map(r.C, r.D)
   {}
@@ -286,7 +286,7 @@
   {}
   #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)
   {}
   #pragma omp target teams map(u.B)  // expected-error {{mapping of union members is not allowed}}
   {}
@@ -434,7 +434,7 @@
   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)
 #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}}
@@ -508,7 +508,7 @@
   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])
 #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}}
diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp
--- a/clang/test/OpenMP/target_update_codegen.cpp
+++ b/clang/test/OpenMP/target_update_codegen.cpp
@@ -291,5 +291,405 @@
   #pragma omp target update from(arg) if(arg) device(4)
   {++arg;}
 }
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp -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-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-64
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp -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-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK5 --check-prefix CK5-32
+
+// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -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-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-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-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-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-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK5
+
+// CK5: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+// CK5-LABEL: lvalue
+void lvalue(int *B, int l, int e) {
+
+  // CK5-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK5-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK5-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK5-DAG: store i32* [[ZERO:%.+]], i32** [[BPC0]]
+  // CK5-DAG: store i32* [[ONE:%.+]], i32** [[PC0]]
+  // CK5-DAG: [[ZERO]] = load i32*, i32** [[B_ADDR:%.+]]
+  // CK5-DAG: [[ONE]] = load i32*, i32** [[B_ADDR]]
+  #pragma omp target update to(*B)
+  *B += e;
+  #pragma omp target update from(*B)
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK6 -verify -fopenmp -fopenmp-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-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-64
+// RUN: %clang_cc1 -DCK6 -verify -fopenmp -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-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK6 --check-prefix CK6-32
+
+// RUN: %clang_cc1 -DCK6 -verify -fopenmp-simd -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-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-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-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-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-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK6
+
+// CK6: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+// CK6-LABEL: lvalue
+void lvalue(int *B, int l, int e) {
+
+  // CK6-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK6-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK6-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK6-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK6-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK6-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK6-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK6-DAG: store i32* [[ZERO:%.+]], i32** [[BPC0]]
+  // CK6-DAG: store i32* [[ADD_PTR:%.+]], i32** [[PC0]]
+  // CK6-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i{{32|64}} [[IDX_EXT:%.+]]
+  // CK6-DAG: [[TWO:%.+]] = load i32, i32* [[L_ADDR:%.+]]
+  // CK6-DAG: store i32 [[L:%.+]], i32* [[L_ADDR]]
+  #pragma omp target update to(*(B+l))
+  *(B+l) += e;
+  #pragma omp target update from(*(B+l))
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK7 -verify -fopenmp -fopenmp-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-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK7 --check-prefix CK7-64
+// RUN: %clang_cc1 -DCK7 -verify -fopenmp -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-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK7 --check-prefix CK7-32
+
+// RUN: %clang_cc1 -DCK7 -verify -fopenmp-simd -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-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-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-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-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-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK7
+
+// CK7: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4]
+// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 33]
+
+// CK7-LABEL: lvalue
+void lvalue(int *B, int l, int e) {
+
+  // CK7-DAG: call void @__tgt_target_data_update(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK7-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK7-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+  // CK7-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK7-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK7-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**
+  // CK7-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32**
+  // CK7-DAG: store i32* [[ZERO:%.+]], i32** [[BPC0]]
+  // CK7-DAG: store i32* [[ARRAY_IDX:%.+]], i32** [[PC0]]
+  // CK7-DAG: [[ZERO]] = load i32*, i32** [[B_ADDR:%.+]]
+  // CK7-DAG: [[ARRAY_IDX]] = getelementptr inbounds i32, i32* [[ADD_PTR:%.+]], i{{32|64}} [[IDX_PROM:%.+]]
+  // CK7-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i64 [[IDX_EXT:%.+]]
+  // CK7-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32, i32* [[ONE:%.+]], i32 [[TWO:%.+]] 
+  // CK7-32-DAG: [[ONE]] = load i32*, i32** [[B_ADDR]]
+  // CK7-32-DAG: [[TWO]] = load i32, i32* [[L_ADDR:%.+]]
+  // CK7-32-DAG: [[IDX_PROM]] = load i32, i32* [[L_ADDR]]
+  // CK7-64-DAG: [[IDX_EXT:%.+]] = sext i32 [[TWO:%.+]] to i64 
+  // CK7-64-DAG: [[TWO:%.+]] = load i32, i32* [[L_ADDR:%.+]]
+  // CK7-64-DAG: [[IDX_PROM]] = sext i32 [[THREE:%.+]] to i64
+  // CK7-64-DAG: [[THREE]] = load i32, i32* [[L_ADDR]]
+  #pragma omp target update to((B+l)[l])
+  (B+l)[l] += e;
+  #pragma omp target update from((B+l)[l])
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp -fopenmp-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-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK8 --check-prefix CK8-64
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp -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-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK8 --check-prefix CK8-32
+
+// RUN: %clang_cc1 -DCK8 -verify -fopenmp-simd -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-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-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-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-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-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-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-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-64
+// RUN: %clang_cc1 -DCK9 -verify -fopenmp -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-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK9 --check-prefix CK9-32
+
+// RUN: %clang_cc1 -DCK9 -verify -fopenmp-simd -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-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-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-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-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-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+
+#ifdef CK9
+
+struct S {
+  double *p;
+};
+
+// CK9: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} 32, i64 281474976710673]
+// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+
+// CK9-LABEL: lvalue
+void lvalue(struct S *s, int l, int e) {
+
+  // CK9-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK9-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK9-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK9-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+  //
+  // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK9-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK9-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
+  // CK9-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
+  // CK9-DAG: store double** [[P:%.+]], double*** [[BPC0]]
+  // CK9-DAG: store double* [[THREE:%.+]], double** [[PC0]]
+  // CK9-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE0]]
+  // CK9-DAG: [[P]] = getelementptr inbounds [[STRUCT_S:%.+]], [[STRUCT_S]]* [[ONE:%.+]], i32 0, i32 0
+  // CK9-DAG: [[ONE]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S_ADDR:%.+]]
+  // CK9-DAG: [[THREE]] = load double*, double** [[P_1:%.+]],
+  // CK9-DAG: [[P_1]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[TWO:%.+]], i32 0, i32 0
+  // CK9-DAG: [[TWO]] = load [[STRUCT_S]]*, [[STRUCT_S]]** [[S_ADDR:%.+]]
+  #pragma omp target update to(*(s->p))
+    *(s->p) += e;
+  #pragma omp target update from(*(s->p))
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK10 -verify -fopenmp -fopenmp-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-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK10 --check-prefix CK10-64
+// RUN: %clang_cc1 -DCK10 -verify -fopenmp -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-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK10 --check-prefix CK10-32
+
+// RUN: %clang_cc1 -DCK10 -verify -fopenmp-simd -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-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-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-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-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-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK10
+
+struct S {
+  double *p;
+};
+
+// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+
+// CK10-LABEL: lvalue
+void lvalue(struct S *s, int l, int e) {
+
+  // CK10-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK10-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK10-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK10-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+  //
+  // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK10-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK10-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
+  // CK10-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
+  // CK10-DAG: store double** [[P_VAL:%.+]], double*** [[BPC0]]
+  // CK10-DAG: store double* [[ADD_PTR:%.+]], double** [[PC0]]
+  // CK10-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE0]]
+  // CK10-64-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[IDX_EXT:%.+]]
+  // CK10-32-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[L_VAL:%.+]]
+  // CK10-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
+  // CK10-DAG: [[THREE]] = load double*, double** [[P_VAL_1:%.+]]
+  // CK10-DAG: getelementptr inbounds {{.+}}, {{.+}}* [[SS_1:%.+]], i32 0, i32 0
+  // CK10-DAG: [[SS_1]] = load {{.+}}*, {{.+}}** [[S_ADDR:%.+]]
+  #pragma omp target update to(*(s->p+l))
+    *(s->p+l) += e;
+  #pragma omp target update from(*(s->p+l))
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK11 -verify -fopenmp -fopenmp-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-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK11 --check-prefix CK11-64
+// RUN: %clang_cc1 -DCK11 -verify -fopenmp -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-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK11 --check-prefix CK11-32
+
+// RUN: %clang_cc1 -DCK11 -verify -fopenmp-simd -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-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-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-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-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-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK11
+
+struct S {
+  double *p;
+};
+// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710674]
+
+// CK11-LABEL: lvalue
+void lvalue(struct S *s, int l, int e) {
+
+  // CK11-DAG: call void @__tgt_target_data_update(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK11-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK11-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK11-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+  //
+  // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK11-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK11-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double***
+  // CK11-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double**
+  // CK11-DAG: store double** [[P:%.+]], double*** [[BPC0]]
+  // CK11-DAG: store double* [[ARRAY_IDX:%.+]], double** [[PC0]]
+  // CK11-DAG: store i{{.+}} 8, i{{.+}} [[SIZE0]]
+  // CK11-DAG: [[P]] = getelementptr inbounds [[STRUCT_S:%.+]], [[STRUCT_S]]* [[SS_1:%.+]], i32 0, i32 0
+  // CK11-DAG: [[ARRAY_IDX]] = getelementptr inbounds double, double* [[ADD_PTR:%.+]], i{{.+}} 3
+  // CK11-64-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[IDX_EXT:%.+]]
+  // CK11-32-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} [[L_VAL:%.+]]
+  // CK11-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64
+  #pragma omp target update to((s->p+l)[3])
+    (s->p+l)[3] += e;
+  #pragma omp target update from((s->p+l)[3])
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK12 -verify -fopenmp -fopenmp-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-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK12 --check-prefix CK12-64
+// RUN: %clang_cc1 -DCK12 -verify -fopenmp -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-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK12 --check-prefix CK12-32
+
+// RUN: %clang_cc1 -DCK12 -verify -fopenmp-simd -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-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-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-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-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-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+#ifdef CK12
+
+struct S {
+  double *p;
+  struct S *sp;
+};
+// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710672, i64 17]
+
+// CK12-LABEL: lvalue
+void lvalue(struct S *s, int l, int e) {
+
+  // CK12-DAG: call void @__tgt_target_data_update(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}})
+  // CK12-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
+  // CK12-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK12-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]]
+  //
+  // CK12-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
+  // CK12-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
+  // CK12-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2
+  // CK12-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to double***
+  // CK12-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double**
+  // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[BPC2]]
+  // CK12-DAG: store double* [[SIX:%.+]], double** [[PC2]]
+  // CK12-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE2]]
+  // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
+  // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
+  // CK12-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1
+  // CK12-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to [[STRUCT_S:%.+]]***
+  // CK12-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double***
+  // CK12-DAG: store [[STRUCT_S]]** [[SP:%.+]], [[STRUCT_S]]*** [[BPC1]]
+  // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[PC1]]
+  // CK12-DAG: store i{{.+}} {{4|8}}, i{{.+}}* [[SIZE1]]
+  // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
+  // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
+  // CK12-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0
+  // CK12-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[STRUCT_S:%.+]]**
+  // CK12-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to [[STRUCT_S]]***
+  // CK12-DAG: store [[STRUCT_S]]** [[SP:%.+]], [[STRUCT_S]]*** [[S_VAL:%.+]]
+  // CK12-DAG: store i{{.+}} {{.+}}, i{{.+}}* [[SIZE0]]
+  // CK12-DAG: [[SP]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[ONE:%.+]], i32 0, i32 1
+  #pragma omp target update to(*(s->sp->p))
+    *(s->sp->p) = e;
+  #pragma omp target update from(*(s->sp->p))
+}
+
 #endif
 #endif
diff --git a/clang/test/OpenMP/target_update_from_messages.cpp b/clang/test/OpenMP/target_update_from_messages.cpp
--- a/clang/test/OpenMP/target_update_from_messages.cpp
+++ b/clang/test/OpenMP/target_update_from_messages.cpp
@@ -94,7 +94,7 @@
 #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)
 #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)
@@ -106,7 +106,7 @@
 
 #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])
 #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}}
@@ -148,7 +148,7 @@
 #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)
 #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)
@@ -160,7 +160,7 @@
 
 #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])
 #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}}
diff --git a/clang/test/OpenMP/target_update_to_messages.cpp b/clang/test/OpenMP/target_update_to_messages.cpp
--- a/clang/test/OpenMP/target_update_to_messages.cpp
+++ b/clang/test/OpenMP/target_update_to_messages.cpp
@@ -94,7 +94,7 @@
 #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)
 #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)
@@ -106,7 +106,7 @@
 
 #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])
 #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}}
@@ -147,7 +147,7 @@
 #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)
 #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)
@@ -159,7 +159,7 @@
 
 #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])
 #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}}