[OpenMP 5.0] target update list items need not be contiguous (Sema)

[cchen]

Summary

Removed a restriction that list items for target update directive must have contiguous storage.

[cchen]

Fix summary message

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

14450

Ad the check for the version of OpenMP standard, it must be allowed only for OpenMP >= 50.

14470–14471

Same here, allow only for OpenMP >= 50. Also, split the patch into 2 parts, one for pointers stuff and another one for target update directive.

[cchen]

Remove the restriction rule for array section to another patch

[cchen]

Rebase and fix add OpenMP 50 guard

[ABataev]

What about codegen tests?

clang/test/OpenMP/target_update_ast_print.cpp

25

It seems to me, the tests are not updated. All non-contiguous tests must pass only if OpenMP version is 50

[cchen]

Add and fix test.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

14470–14472

SemaRef.getLangOpts().OpenMP < 50 || DKind != OMPD_target_update

clang/test/OpenMP/target_update_codegen.cpp

325–326 ↗

(On Diff #229922)

We still map the whole array here, right?

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_ordered_directive_with_param : Error<
   "'ordered' directive %select{without any clauses|with 'threads' clause}0 cannot be closely nested inside ordered region with specified parameter">;
 def err_omp_ordered_directive_without_param : Error<
   "'ordered' directive with 'depend' clause cannot be closely nested inside ordered region without specified parameter">;
 def note_omp_ordered_param : Note<
   "'ordered' clause with specified parameter">;
 def err_omp_expected_base_var_name : Error<
   "expected variable name as a base of the array %select{subscript|section}0">;
+def err_omp_pointer_type_not_last : Error<
+  "only one level of indirection allowed in array section">;
 def err_omp_map_shared_storage : Error<
   "variable already marked as mapped in current construct">;
 def err_omp_invalid_map_type_for_directive : Error<
   "%select{map type '%1' is not allowed|map type must be specified}0 for '#pragma omp %2'">;
 def err_omp_no_clause_for_directive : Error<
   "expected at least one %0 clause for '#pragma omp %1'">;
 def err_omp_threadprivate_in_clause : Error<
   "threadprivate variables are not allowed in '%0' clause">;

clang/lib/Sema/SemaOpenMP.cpp

 
 //===----------------------------------------------------------------------===//
 // Stack of data-sharing attributes for variables
 //===----------------------------------------------------------------------===//
 
 static const Expr *checkMapClauseExpressionBase(
     Sema &SemaRef, Expr *E,
     OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
-    OpenMPClauseKind CKind, bool NoDiagnose);
+    OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose);
 
 namespace {
 /// Default data sharing attributes, which can be applied to directive.
 enum DefaultDataSharingAttributes {
   DSA_unspecified = 0, /// Data sharing attribute not specified.
   DSA_none = 1 << 0,   /// Default data sharing attribute 'none'.
   DSA_shared = 1 << 1, /// Default data sharing attribute 'shared'.
 };
if (auto *TE = dyn_cast<CXXThisExpr>(E->getBase()->IgnoreParens())) {
         if (DVar.CKind != OMPC_unknown)
           ImplicitFirstprivate.push_back(E);
       }
       return;
     }
     if (isOpenMPTargetExecutionDirective(DKind)) {
       OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
       if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map,
-                                        /*NoDiagnose=*/true))
+                                        Stack->getCurrentDirective(), /*NoDiagnose=*/true))
         return;
       const auto *VD = cast<ValueDecl>(
           CurComponents.back().getAssociatedDeclaration()->getCanonicalDecl());
       if (!Stack->checkMappableExprComponentListsForDecl(
               VD, /*CurrentRegionOnly=*/true,
               [&CurComponents](
                   OMPClauseMappableExprCommon::MappableExprComponentListRef
                       StackComponents,

 
 // 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) {
+    OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose) {
   SourceLocation ELoc = E->getExprLoc();
   SourceRange ERange = E->getSourceRange();
 
   // The base of elements of list in a map clause have to be either:
   //  - a reference to variable or field.
   //  - a member expression.
   //  - an array expression.
   //
   // E.g. if we have the expression 'r.S.Arr[:12]', we want to retrieve the
   // reference to 'r'.
   //
   // If we have:
   //
   // struct SS {
   //   Bla S;
   //   foo() {
   //     #pragma omp target map (S.Arr[:12]);
   //   }
   // }
   //
   // We want to retrieve the member expression 'this->S';
 
   const Expr *RelevantExpr = nullptr;
 
+  // OpenMP 5.0
+  // The target update construct was modified to allow array section that
+  // specify discontiguous storage.
+  //
   // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.2]
   //  If a list item is an array section, it must specify contiguous storage.
   //
   // For this restriction it is sufficient that we make sure only references
   // to variables or fields and array expressions, and that no array sections
   // exist except in the rightmost expression (unless they cover the whole
   // dimension of the array). E.g. these would be invalid:
   //
   //   r.ArrS[3:5].Arr[6:7]
   //
   //   r.ArrS[3:5].x
   //
   // but these would be valid:
   //   r.ArrS[3].Arr[6:7]
   //
   //   r.ArrS[3].x
 
   bool AllowUnitySizeArraySection = true;
   bool AllowWholeSizeArraySection = true;
 
+  // OpenMP 5.0 [2.1.5, Array Sections]
+  // If a consecutive sequence of array subscript expressions appears in an
+  // array section, and the first subscript expression in the sequence uses
+  // the extended array section syntax defined in this section, then only the
+  // last subscript expression in the sequence may select array elements that
+  // have a pointer type.
+  bool IsLastArraySection = true;
+  bool HasMeetPointer = false;
+
   while (!RelevantExpr) {
     E = E->IgnoreParenImpCasts();
 
     if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
       if (!isa<VarDecl>(CurE->getDecl()))
         return nullptr;
 
       RelevantExpr = CurE;
if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
       bool IsPointer = CurType->isAnyPointerType();
 
       if (!IsPointer && !CurType->isArrayType()) {
         SemaRef.Diag(ELoc, diag::err_omp_expected_base_var_name)
             << 0 << CurE->getSourceRange();
         return nullptr;
       }
 
+      if (HasMeetPointer && !IsLastArraySection) {
+        SemaRef.Diag(ELoc, diag::err_omp_pointer_type_not_last)
+            << CurE->getSourceRange();
+        return nullptr;
+      }
+
       bool NotWhole =
           checkArrayExpressionDoesNotReferToWholeSize(SemaRef, CurE, CurType);
       bool NotUnity =
           checkArrayExpressionDoesNotReferToUnitySize(SemaRef, CurE, CurType);
 
       if (AllowWholeSizeArraySection) {
         // Any array section is currently allowed. Allowing a whole size array
         // section implies allowing a unity array section as well.
         //
         // If this array section refers to the whole dimension we can still
         // accept other array sections before this one, except if the base is a
         // pointer. Otherwise, only unitary sections are accepted.
         if (NotWhole || IsPointer)
           AllowWholeSizeArraySection = false;
-      } else if (AllowUnitySizeArraySection && NotUnity) {
+      } else if (DKind != OMPD_target_update &&
+         (AllowUnitySizeArraySection && NotUnity)) {
         // A unity or whole array section is not allowed and that is not
         // compatible with the properties of the current array section.
         SemaRef.Diag(
             ELoc, diag::err_array_section_does_not_specify_contiguous_storage)
             << CurE->getSourceRange();
         return nullptr;
       }
 
       if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
if (auto *CurE = dyn_cast<DeclRefExpr>(E)) {
                          diag::err_omp_invalid_map_this_expr);
             SemaRef.Diag(CurE->getLowerBound()->getExprLoc(),
                          diag::note_omp_invalid_lower_bound_on_this_ptr_mapping);
           }
         }
         RelevantExpr = TE;
       }
 
+      IsLastArraySection = false;
+      if (IsPointer) HasMeetPointer = true;
+
       // Record the component - we don't have any declaration associated.
       CurComponents.emplace_back(CurE, nullptr);
     } else {
       if (!NoDiagnose) {
         // If nothing else worked, this is not a valid map clause expression.
         SemaRef.Diag(
             ELoc, diag::err_omp_expected_named_var_member_or_array_expression)
             << ERange;
for (Expr *RE : MVLI.VarList) {
     }
 
     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.
     const Expr *BE = checkMapClauseExpressionBase(
-        SemaRef, SimpleExpr, CurComponents, CKind, /*NoDiagnose=*/false);
+        SemaRef, SimpleExpr, CurComponents, CKind, DSAS->getCurrentDirective(),
+         /*NoDiagnose=*/false);
     if (!BE)
       continue;
 
     assert(!CurComponents.empty() &&
            "Invalid mappable expression information.");
 
     if (const auto *TE = dyn_cast<CXXThisExpr>(BE)) {
       // Add store "this" pointer to class in DSAStackTy for future checking

clang/test/OpenMP/target_map_codegen.cpp

   // CK19-DAG: [[SEC2]] = getelementptr {{.*}}double* [[SEC22:[^,]+]], i{{.+}} 0
   // CK19-DAG: [[SEC22]] = load double*, double** [[SEC222:%[^,]+]],
   // CK19-DAG: [[SEC222]] = getelementptr {{.*}}double** [[SEC2222:[^,]+]], i{{.+}} 2
   // CK19-DAG: [[SEC2222]] = load double**, double*** [[SEC22222:%[^,]+]],
   // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}double*** [[SEC222222:[^,]+]], i{{.+}} 0
   // CK19-DAG: [[SEC222222]] = load double***, double**** [[PTR]],
 
   // CK19: call void [[CALL42:@.+]](double*** {{[^,]+}})
-  #pragma omp target map(mptras[:1][2][:13])
+  #pragma omp target map(mptras[0][2][:13])
   {
     mptras[1][2][3]++;
   }
 
   // Region 43 - the memory is not contiguous for this map - will map the whole last dimension.
   // CK19-DAG: call i32 @__tgt_target(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE43]]{{.+}})
   // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
   // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]

clang/test/OpenMP/target_map_messages.cpp

void SAclient(int arg) {
   {}
   #pragma omp target map(mvla[:1][:][:])
   {}
   #pragma omp target map(mvla2[:1][:2][:11])
   {}
   #pragma omp target map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}}
   {}
 
-  #pragma omp target map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{array section does not specify contiguous storage}}
+  #pragma omp target map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{only one level of indirection allowed in array section}}
   {}
-  #pragma omp target map(mptr[:1][:2-1][2:4-3])
+  #pragma omp target map(mptr[:1][:2-1][2:4-3]) // expected-error {{only one level of indirection allowed in array section}}
   {}
-  #pragma omp target map(mptr[:1][:arg][2:4-3]) // correct if arg is 1.
+  #pragma omp target map(mptr[:1][:arg][2:4-3]) // expected-error {{only one level of indirection allowed in array section}}
   {}
-  #pragma omp target map(mptr[:1][:2-1][0:2])
+  #pragma omp target map(mptr[:1][:2-1][0:2]) // expected-error {{only one level of indirection allowed in array section}}
   {}
-  #pragma omp target map(mptr[:1][:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
+  #pragma omp target map(mptr[:1][:2][0:2]) // expected-error {{only one level of indirection allowed in array section}}
   {}
   #pragma omp target map(mptr[:1][:][0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
   {}
-  #pragma omp target map(mptr[:2][:1][0:2]) // expected-error {{array section does not specify contiguous storage}}
+  #pragma omp target map(mptr[:2][:1][0:2]) // expected-error {{only one level of indirection allowed in array section}}
   {}
 
   #pragma omp target map(r.ArrS[0].B)
   {}
   #pragma omp target map(r.ArrS[:1].B) // expected-error {{OpenMP array section is not allowed here}}
   {}
   #pragma omp target map(r.ArrS[:arg].B) // expected-error {{OpenMP array section is not allowed here}}
   {}

clang/test/OpenMP/target_teams_map_messages.cpp

void SAclient(int arg) {
   {}
   #pragma omp target teams map(mvla[:1][:][:])
   {}
   #pragma omp target teams map(mvla2[:1][:2][:11])
   {}
   #pragma omp target teams map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}}
   {}
 
-  #pragma omp target teams map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{array section does not specify contiguous storage}}
+  #pragma omp target teams map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{only one level of indirection allowed in array section}}
   {}
-  #pragma omp target teams map(mptr[:1][:2-1][2:4-3])
+  #pragma omp target teams map(mptr[:1][:2-1][2:4-3]) // expected-error {{only one level of indirection allowed in array section}}
   {}
-  #pragma omp target teams map(mptr[:1][:arg][2:4-3]) // correct if arg is 1.
+  #pragma omp target teams map(mptr[:1][:arg][2:4-3]) // expected-error {{only one level of indirection allowed in array section}}
   {}
-  #pragma omp target teams map(mptr[:1][:2-1][0:2])
+  #pragma omp target teams map(mptr[:1][:2-1][0:2]) // expected-error {{only one level of indirection allowed in array section}}
   {}
-  #pragma omp target teams map(mptr[:1][:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
+  #pragma omp target teams map(mptr[:1][:2][0:2]) // expected-error {{only one level of indirection allowed in array section}}
   {}
   #pragma omp target teams map(mptr[:1][:][0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
   {}
-  #pragma omp target teams map(mptr[:2][:1][0:2]) // expected-error {{array section does not specify contiguous storage}}
+  #pragma omp target teams map(mptr[:2][:1][0:2]) // expected-error {{only one level of indirection allowed in array section}}
   {}
 
   #pragma omp target teams map(r.ArrS[0].B)
   {}
   #pragma omp target teams map(r.ArrS[:1].B) // expected-error {{OpenMP array section is not allowed here}}
   {}
   #pragma omp target teams map(r.ArrS[:arg].B) // expected-error {{OpenMP array section is not allowed here}}
   {}

clang/test/OpenMP/target_update_ast_print.cpp

 template <class T, class U>
 T foo(T targ, U uarg) {
   static T a;
   U b;
   int l;
 #pragma omp target update to(a) if(l>5) device(l) nowait depend(inout:l)
 
 #pragma omp target update from(b) if(l<5) device(l-1) nowait depend(inout:l)
+
+  U marr[10][10][10];
+#pragma omp target update to(marr[2][0:2][0:2])
+
+#pragma omp target update from(marr[2][0:2][0:2])
+
+#pragma omp target update to(marr[:][0:2][0:2])
+
+#pragma omp target update from(marr[:][0:2][0:2])
+
+#pragma omp target update to(marr[:][:l][l:])
+
+#pragma omp target update from(marr[:][:l][l:])
+
+#pragma omp target update to(marr[:2][:1][:])
+
+#pragma omp target update from(marr[:2][:1][:])
+
+#pragma omp target update to(marr[:2][:][:1])
+
+#pragma omp target update from(marr[:2][:][:1])
+
+#pragma omp target update to(marr[:2][:][1:])
+
+#pragma omp target update from(marr[:2][:][1:])
+
+#pragma omp target update to(marr[:1][3:2][:2])
+
+#pragma omp target update from(marr[:1][3:2][:2])
+
+#pragma omp target update to(marr[:1][:2][0])
+
+#pragma omp target update from(marr[:1][:2][0])
+
   return a + targ + (T)b;
 }
 // CHECK:      static T a;
 // CHECK-NEXT: U b;
 // CHECK-NEXT: int l;
 // CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) nowait depend(inout : l){{$}}
 // CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) nowait depend(inout : l)
 // CHECK:      static int a;
 // CHECK-NEXT: float b;
 // CHECK-NEXT: int l;
 // CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) nowait depend(inout : l)
 // CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) nowait depend(inout : l)
 // CHECK:      static char a;
 // CHECK-NEXT: float b;
 // CHECK-NEXT: int l;
 // CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) nowait depend(inout : l)
 // CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) nowait depend(inout : l)
+// CHECK-NEXT: marr[10][10][10];
+// CHECK-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
+// CHECK-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
+// CHECK-NEXT: #pragma omp target update to(marr[:][0:2][0:2])
+// CHECK-NEXT: #pragma omp target update from(marr[:][0:2][0:2])
+// CHECK-NEXT: #pragma omp target update to(marr[:][:l][l:])
+// CHECK-NEXT: #pragma omp target update from(marr[:][:l][l:])
+// CHECK-NEXT: #pragma omp target update to(marr[:2][:1][:])
+// CHECK-NEXT: #pragma omp target update from(marr[:2][:1][:])
+// CHECK-NEXT: #pragma omp target update to(marr[:2][:][:1])
+// CHECK-NEXT: #pragma omp target update from(marr[:2][:][:1])
+// CHECK-NEXT: #pragma omp target update to(marr[:2][:][1:])
+// CHECK-NEXT: #pragma omp target update from(marr[:2][:][1:])
+// CHECK-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
+// CHECK-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
+// CHECK-NEXT: #pragma omp target update to(marr[:1][:2][0])
+// CHECK-NEXT: #pragma omp target update from(marr[:1][:2][0])
 
 int main(int argc, char **argv) {
   static int a;
   int n;
   float f;
 
 // CHECK:      static int a;
 // CHECK-NEXT: int n;
 // CHECK-NEXT: float f;
 #pragma omp target update to(a) if(f>0.0) device(n) nowait depend(in:n)
 // CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n) nowait depend(in : n)
 #pragma omp target update from(f) if(f<0.0) device(n+1) nowait depend(in:n)
 // CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1) nowait depend(in : n)
+
+  float marr[10][10][10];
+// CHECK-NEXT: marr[10][10][10];
+#pragma omp target update to(marr[2][0:2][0:2])
+// CHECK-NEXT: #pragma omp target update to(marr[2][0:2][0:2])
+#pragma omp target update from(marr[2][0:2][0:2])
+// CHECK-NEXT: #pragma omp target update from(marr[2][0:2][0:2])
+#pragma omp target update to(marr[:][0:2][0:2])
+// CHECK-NEXT: #pragma omp target update to(marr[:][0:2][0:2])
+#pragma omp target update from(marr[:][0:2][0:2])
+// CHECK-NEXT: #pragma omp target update from(marr[:][0:2][0:2])
+#pragma omp target update to(marr[:][:n][n:])
+// CHECK-NEXT: #pragma omp target update to(marr[:][:n][n:])
+#pragma omp target update from(marr[:2][:1][:])
+// CHECK-NEXT: #pragma omp target update from(marr[:2][:1][:])
+#pragma omp target update to(marr[:2][:][:1])
+// CHECK-NEXT: #pragma omp target update to(marr[:2][:][:1])
+#pragma omp target update from(marr[:2][:][:1])
+// CHECK-NEXT: #pragma omp target update from(marr[:2][:][:1])
+#pragma omp target update to(marr[:2][:][1:])
+// CHECK-NEXT: #pragma omp target update to(marr[:2][:][1:])
+#pragma omp target update from(marr[:2][:][1:])
+// CHECK-NEXT: #pragma omp target update from(marr[:2][:][1:])
+#pragma omp target update to(marr[:1][3:2][:2])
+// CHECK-NEXT: #pragma omp target update to(marr[:1][3:2][:2])
+#pragma omp target update from(marr[:1][3:2][:2])
+// CHECK-NEXT: #pragma omp target update from(marr[:1][3:2][:2])
+#pragma omp target update to(marr[:1][:2][0])
+// CHECK-NEXT: #pragma omp target update to(marr[:1][:2][0])
+#pragma omp target update from(marr[:1][:2][0])
+// CHECK-NEXT: #pragma omp target update from(marr[:1][:2][0])
   return foo(argc, f) + foo(argv[0][0], f) + a;
 }
 
 #endif

clang/test/OpenMP/target_update_from_messages.cpp

T tmain(T argc) {
   T &j = i;
   T *k = &j;
   T x;
   T y;
   T from;
   const T (&l)[5] = da;
   T *m;
   S7 s7;
+  T ***mptr;
 
 #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(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update from(s7.p[:10])
 #pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}}
 #pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
 #pragma omp target data map(to: s7.i)
   {
 #pragma omp target update from(s7.x)
   }
+#pragma omp target update from(mptr[:2][2+2-4:1][0:5+5]) // expected-error 2 {{only one level of indirection allowed in array section}}
+#pragma omp target update from(mptr[:1][:2-1][2:4-3]) // expected-error 2 {{only one level of indirection allowed in array section}}
+#pragma omp target update from(mptr[:1][:2][0:2]) // expected-error 2 {{only one level of indirection allowed in array section}}
 
   return 0;
 }
 
 int main(int argc, char **argv) {
   const int d = 5;
   const int da[5] = { 0 };
   S4 e(4);
   S5 g(5);
   int i, t[20];
   int &j = i;
   int *k = &j;
   int x;
   int y;
   int from;
   const int (&l)[5] = da;
   int *m;
   S7 s7;
+  int ***mptr;
 
 #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(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}}
 #pragma omp target update from(s7.p[:10])
 #pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}}
 #pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
 #pragma omp target data map(to: s7.i)
   {
 #pragma omp target update from(s7.x)
   }
+#pragma omp target update from(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(mptr[:1][:2-1][2:4-3]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(mptr[:1][:2][0:2]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 
   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

T tmain(T argc) {
   T i, t[20];
   T &j = i;
   T *k = &j;
   T x;
   T y;
   T to;
   const T (&l)[5] = da;
   S7 s7;
+  T ***mptr;
 
 #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(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(mptr[:2][2+2-4:1][0:5+5]) // expected-error 2 {{only one level of indirection allowed in array section}}
+#pragma omp target update to(mptr[:1][:2-1][2:4-3]) // expected-error 2 {{only one level of indirection allowed in array section}}
+#pragma omp target update to(mptr[:1][:2][0:2]) // expected-error 2 {{only one level of indirection allowed in array section}}
   return 0;
 }
 
 int main(int argc, char **argv) {
   const int d = 5;
   const int da[5] = { 0 };
   S4 e(4);
   S5 g(5);
   int i, t[20];
   int &j = i;
   int *k = &j;
   int x;
   int y;
   int to;
   const int (&l)[5] = da;
   S7 s7;
   int *m;
+  int ***mptr;
 
 #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(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(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(mptr[:1][:2-1][2:4-3]) // expected-error {{only one level of indirection allowed in array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(mptr[:1][:2][0:2]) // expected-error {{only one level of indirection allowed in array section}} 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}}
 }