[OpenMP5.0] map item can be non-contiguous for target update

[cchen]

Summary

In order not to modify the tgt_target_data_update information but still be
able to pass the extra information for non-contiguous map item (offset,
count, and stride for each dimension), this patch overload arg when
the maptype is set as OMP_MAP_DESCRIPTOR. The origin arg is for
passing the pointer information, however, the overloaded arg is an
array of descriptor_dim:

struct descriptor_dim {

int64_t offset;
int64_t count;
int64_t stride

};

and the array size is the same as dimension size. In addition, since we
have count and stride information in descriptor_dim, we can replace/overload the
arg_size parameter by using dimension size.

For supporting stride in array section, we use a dummy dimension in
descriptor to store the unit size. The formula for counting the stride
in dimension D_n: unit size * (D_0 * D_1 ... * D_n-1) * D_n.stride.

Demonstrate how it works:

double arr[3][4][5];

D0: { offset = 0, count = 1, stride = 8 }                                // offset, count, dimension size always be 0, 1, 1 for this extra dimension, stride is the unit size
D1: { offset = 0, count = 2, stride = 8 * 1 * 2 = 16 }                   // stride = unit size * (product of dimension size of D0) * D1.stride = 4 * 1 * 2 = 8
D2: { offset = 2, count = 2, stride = 8 * (1 * 5) * 1 = 40  }            // stride = unit size * (product of dimension size of D0, D1) * D2.stride = 4 * 5 * 1 = 20
D3: { offset = 0, count = 2, stride = 8 * (1 * 5 * 4) * 2 = 320 }        // stride = unit size * (product of dimension size of D0, D1, D2) * D3.stride = 4 * 25 * 2 = 200

// X here means we need to offload this data, therefore, runtime will transfer
// data from offset 80, 96, 120, 136, 400, 416, 440, 456
// Runtime patch: https://reviews.llvm.org/D82245
// OOOOO OOOOO OOOOO
// OOOOO OOOOO OOOOO
// XOXOO OOOOO XOXOO
// XOXOO OOOOO XOXOO

[ABataev]

Also, what about compatibility with declare mapper? Can you add tests for it?

clang/lib/CodeGen/CGOpenMPRuntime.cpp

7877–7881

DO not use braced initializer list https://llvm.org/docs/CodingStandards.html#id26

7907–7915

No need for braces here per coding standard

7944

have already have

7960

debug code

7961

Move this after the first if statement

7962–7964

Better to make it something like:

if (const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr))

8012–8013

If no stride at all, why need to divide?

8035

Same, if no stride at all, no need to mul.

clang/lib/Sema/SemaOpenMP.cpp

17129–17135

Better to transform it to something like:

if (!OASE || OASE->getLength())
  break;
SemaRef.Diag(...)

[cchen]

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

[cchen]

Fix based on feedback

[ABataev]

In D84192#2165235, @cchen wrote:

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

[cchen]

In D84192#2165249, @ABataev wrote:

In D84192#2165235, @cchen wrote:

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

Oh, I didn't see that. I'll add test for it, thanks!

[cchen]

In D84192#2165249, @ABataev wrote:

In D84192#2165235, @cchen wrote:

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

[ABataev]

In D84192#2165396, @cchen wrote:

In D84192#2165249, @ABataev wrote:

In D84192#2165235, @cchen wrote:

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

[cchen]

In D84192#2165438, @ABataev wrote:

In D84192#2165396, @cchen wrote:

In D84192#2165249, @ABataev wrote:

In D84192#2165235, @cchen wrote:

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

[ABataev]

In D84192#2165517, @cchen wrote:

In D84192#2165438, @ABataev wrote:

In D84192#2165396, @cchen wrote:

In D84192#2165249, @ABataev wrote:

In D84192#2165235, @cchen wrote:

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

And it is fine. How does it work in declare mapper, the main question? Does it support extended array sections format mapoers with maps, to and from? Shall it be supported in declare mapper at all?

[cchen]

In D84192#2165522, @ABataev wrote:

In D84192#2165517, @cchen wrote:

In D84192#2165438, @ABataev wrote:

In D84192#2165396, @cchen wrote:

In D84192#2165249, @ABataev wrote:

In D84192#2165235, @cchen wrote:

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

And it is fine. How does it work in declare mapper, the main question? Does it support extended array sections format mapoers with maps, to and from? Shall it be supported in declare mapper at all?

After discussing with @dreachem, my understanding is that it is not incorrect to not allowing extended/non-contiguous array section to appear in declare mapper.
For the declare mapper directive, since spec only allows map clause, extended array section (with stride) or non-contiguous array section are both not allowed.
For using the mapper in map/to/from clause, if mapping or updating an array section of type T, where there is a mapper declared for T, then the mapper needs to apply as if to each element of the array or array section. Spec is intentionally not sufficiently clear on this for target update so the semantic check in Clang is not incorrect. Which lead to the fact that I might not need to support extended/non-contiguous array section for declare mapper.

[ABataev]

In D84192#2167636, @cchen wrote:

In D84192#2165522, @ABataev wrote:

In D84192#2165517, @cchen wrote:

In D84192#2165438, @ABataev wrote:

In D84192#2165396, @cchen wrote:

In D84192#2165249, @ABataev wrote:

In D84192#2165235, @cchen wrote:

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

And it is fine. How does it work in declare mapper, the main question? Does it support extended array sections format mapoers with maps, to and from? Shall it be supported in declare mapper at all?

After discussing with @dreachem, my understanding is that it is not incorrect to not allowing extended/non-contiguous array section to appear in declare mapper.
For the declare mapper directive, since spec only allows map clause, extended array section (with stride) or non-contiguous array section are both not allowed.
For using the mapper in map/to/from clause, if mapping or updating an array section of type T, where there is a mapper declared for T, then the mapper needs to apply as if to each element of the array or array section. Spec is intentionally not sufficiently clear on this for target update so the semantic check in Clang is not incorrect. Which lead to the fact that I might not need to support extended/non-contiguous array section for declare mapper.

What exactly is incorrect in clang sema?

[cchen]

In D84192#2167649, @ABataev wrote:

In D84192#2167636, @cchen wrote:

In D84192#2165522, @ABataev wrote:

In D84192#2165517, @cchen wrote:

In D84192#2165438, @ABataev wrote:

In D84192#2165396, @cchen wrote:

In D84192#2165249, @ABataev wrote:

In D84192#2165235, @cchen wrote:

In D84192#2164885, @ABataev wrote:

Also, what about compatibility with declare mapper? Can you add tests for it?

There's a restriction for map clause that non-contiguous is not allowed and I guess it also applies to declare mapper.

I see that to/from clauses allow to use mappers too:

to([mapper(mapper-identifier):]locator-list) from([mapper(mapper-identifier):]locator-list

I'm confused of how to use to([mapper(mapper-identifier):]locator-list) with array section.
For this case:

class C {
public:
  int a;
  double b[3][4][5];
};

#pragma omp declare mapper(id: C s) map(s.a, s.b[0:3][0:4][0:5])

#pragma omp target update from(mapper(id): c)

Clang already has a semantic check in from clause when mapper is used: "mapper type must be of struct, union or class type".
So the only item I can put in from clause in the above example is c and I cannot put c.b[0:2][1:2][0:2] or any even c.b[0:2].

Does clang compile your example? If not, shall it be allowed for to/from clauses or not?

Clang can compile my example but the problem is that Clang do not accept something like #pragma omp target update from(mapper(id): c.b[0:2][1:2][2:2]) (non-contiguous) or even #pragma omp target update from(mapper(id): c.b[2][1][0:2]) (contiguous).
Actually, Clang now only accepts c as struct/union/class type in from(mapper(id): c). And the reason for the restriction is from declare mapper directive - "The type must be of struct, union or class type in C and C++".

And it is fine. How does it work in declare mapper, the main question? Does it support extended array sections format mapoers with maps, to and from? Shall it be supported in declare mapper at all?

After discussing with @dreachem, my understanding is that it is not incorrect to not allowing extended/non-contiguous array section to appear in declare mapper.
For the declare mapper directive, since spec only allows map clause, extended array section (with stride) or non-contiguous array section are both not allowed.
For using the mapper in map/to/from clause, if mapping or updating an array section of type T, where there is a mapper declared for T, then the mapper needs to apply as if to each element of the array or array section. Spec is intentionally not sufficiently clear on this for target update so the semantic check in Clang is not incorrect. Which lead to the fact that I might not need to support extended/non-contiguous array section for declare mapper.

What exactly is incorrect in clang sema?

I mean "not" incorrect for not allowing extended/non-contiguous array section in `to/from([mapper(mapper-identifier):] location-list).

[cchen]

@ABataev, is there any other concern for this patch?

[ABataev]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8960

1. The second argument must be of integer type, not boolean.
2. Why it is signed?

clang/lib/Sema/SemaOpenMP.cpp

17131

No need for else here

[cchen]

Fix coding style and argument on getIntTypeForBitwidth

[ABataev]

LG

[cchen]

@ABataev, do I need to wait for the runtime patch to commit this one? If so, do you have some recommend reviewers for me to add to that patch? I have pinged several times for that patch but haven't got many reviews for it. Thanks!

[ABataev]

In D84192#2172482, @cchen wrote:

@ABataev, do I need to wait for the runtime patch to commit this one? If so, do you have some recommend reviewers for me to add to that patch? I have pinged several times for that patch but haven't got many reviews for it. Thanks!

Yes, it must be committed after runtime patch.

[cchen]

Rebase

[MaskRay]

map item can be non-contiguous for target update

The subject appears to describe a situation, but patch subjects should usually be imperative

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8572

Did you forget a pair of braces?

Please verify whether f2e479db92452594dfe6aff45ab7841f0a8c69a5 is correct.

[cchen]

clang/lib/CodeGen/CGOpenMPRuntime.cpp

8572

Yes, I forgot a pair of braces while rebase and can confirm that f2e479db92452594dfe6aff45ab7841f0a8c69a5 is correct. Thanks for your help!

clang/include/clang/AST/OpenMPClause.h

 #include "clang/AST/NestedNameSpecifier.h"
 #include "clang/AST/Stmt.h"
 #include "clang/AST/StmtIterator.h"
 #include "clang/Basic/LLVM.h"
 #include "clang/Basic/OpenMPKinds.h"
 #include "clang/Basic/SourceLocation.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/MapVector.h"
+#include "llvm/ADT/PointerIntPair.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/iterator.h"
 #include "llvm/ADT/iterator_range.h"
 #include "llvm/Frontend/OpenMP/OMPConstants.h"
 #include "llvm/Frontend/OpenMP/OMPContext.h"
 #include "llvm/Support/Casting.h"
 #include "llvm/Support/Compiler.h"
 #include "llvm/Support/TrailingObjects.h"

 public:
   /// Class that represents a component of a mappable expression. E.g.
   /// for an expression S.a, the first component is a declaration reference
   /// expression associated with 'S' and the second is a member expression
   /// associated with the field declaration 'a'. If the expression is an array
   /// subscript it may not have any associated declaration. In that case the
   /// associated declaration is set to nullptr.
   class MappableComponent {
-    /// Expression associated with the component.
-    Expr *AssociatedExpression = nullptr;
+    /// Pair of Expression and Non-contiguous pair  associated with the
+    /// component.
+    llvm::PointerIntPair<Expr *, 1, bool> AssociatedExpressionNonContiguousPr;
 
     /// Declaration associated with the declaration. If the component does
     /// not have a declaration (e.g. array subscripts or section), this is set
     /// to nullptr.
     ValueDecl *AssociatedDeclaration = nullptr;
 
   public:
     explicit MappableComponent() = default;
     explicit MappableComponent(Expr *AssociatedExpression,
-                               ValueDecl *AssociatedDeclaration)
-        : AssociatedExpression(AssociatedExpression),
+                               ValueDecl *AssociatedDeclaration,
+                               bool IsNonContiguous)
+        : AssociatedExpressionNonContiguousPr(AssociatedExpression,
+                                              IsNonContiguous),
           AssociatedDeclaration(
               AssociatedDeclaration
                   ? cast<ValueDecl>(AssociatedDeclaration->getCanonicalDecl())
                   : nullptr) {}
 
-    Expr *getAssociatedExpression() const { return AssociatedExpression; }
+    Expr *getAssociatedExpression() const {
+      return AssociatedExpressionNonContiguousPr.getPointer();
+    }
+
+    bool isNonContiguous() const {
+      return AssociatedExpressionNonContiguousPr.getInt();
+    }
 
     ValueDecl *getAssociatedDeclaration() const {
       return AssociatedDeclaration;
     }
   };
 
   // List of components of an expression. This first one is the whole
   // expression and the last one is the base expression.

clang/include/clang/Basic/DiagnosticSemaKinds.td

 def err_omp_expected_addressable_lvalue_or_array_item : Error<
   "expected addressable lvalue expression, array element%select{ or array section|, array section or array shaping expression}0%select{| of non 'omp_depend_t' type}1">;
 def err_omp_expected_named_var_member_or_array_expression: Error<
   "expected expression containing only member accesses and/or array sections based on named variables">;
 def err_omp_bit_fields_forbidden_in_clause : Error<
   "bit fields cannot be used to specify storage in a '%0' clause">;
 def err_array_section_does_not_specify_contiguous_storage : Error<
   "array section does not specify contiguous storage">;
+def err_array_section_does_not_specify_length : Error<
+  "array section does not specify length for outermost dimension">;
 def err_omp_union_type_not_allowed : Error<
   "mapping of union members is not allowed">;
 def err_omp_expected_access_to_data_field : Error<
   "expected access to data field">;
 def err_omp_multiple_array_items_in_map_clause : Error<
   "multiple array elements associated with the same variable are not allowed in map clauses of the same construct">;
 def err_omp_duplicate_map_type_modifier : Error<
   "same map type modifier has been specified more than once">;

clang/lib/CodeGen/CGOpenMPRuntime.cpp

enum OpenMPOffloadMappingFlags : uint64_t {
     /// Implicit map
     OMP_MAP_IMPLICIT = 0x200,
     /// Close is a hint to the runtime to allocate memory close to
     /// the target device.
     OMP_MAP_CLOSE = 0x400,
     /// 0x800 is reserved for compatibility with XLC.
     /// Produce a runtime error if the data is not already allocated.
     OMP_MAP_PRESENT = 0x1000,
+    /// Signal that the runtime library should use args as an array of
+    /// descriptor_dim pointers and use args_size as dims. Used when we have
+    /// non-contiguous list items in target update directive
+    OMP_MAP_NON_CONTIG = 0x100000000000,
     /// The 16 MSBs of the flags indicate whether the entry is member of some
     /// struct/class.
     OMP_MAP_MEMBER_OF = 0xffff000000000000,
     LLVM_MARK_AS_BITMASK_ENUM(/* LargestFlag = */ OMP_MAP_MEMBER_OF),
   };
 
   /// Get the offset of the OMP_MAP_MEMBER_OF field.
   static unsigned getFlagMemberOffset() {
public:
     const ValueDecl *getDevicePtrDecl() const { return DevPtrDecl; }
     void setDevicePtrDecl(const ValueDecl *D) { DevPtrDecl = D; }
   };
 
   using MapBaseValuesArrayTy = SmallVector<BasePointerInfo, 4>;
   using MapValuesArrayTy = SmallVector<llvm::Value *, 4>;
   using MapFlagsArrayTy = SmallVector<OpenMPOffloadMappingFlags, 4>;
   using MapMappersArrayTy = SmallVector<const ValueDecl *, 4>;
+  using MapDimArrayTy = SmallVector<uint64_t, 4>;
+  using MapNonContiguousArrayTy = SmallVector<MapValuesArrayTy, 4>;
 
   /// This structure contains combined information generated for mappable
-  /// clauses, including base pointers, pointers, sizes, map types, and
-  /// user-defined mappers.
+  /// clauses, including base pointers, pointers, sizes, map types, user-defined
+  /// mappers, and non-contiguous information.
   struct MapCombinedInfoTy {
+    struct StructNonContiguousInfo {
+      bool IsNonContiguous = false;
+      MapDimArrayTy Dims;
+      MapNonContiguousArrayTy Offsets;
+      MapNonContiguousArrayTy Counts;
+      MapNonContiguousArrayTy Strides;
+    };
     MapBaseValuesArrayTy BasePointers;
     MapValuesArrayTy Pointers;
     MapValuesArrayTy Sizes;
     MapFlagsArrayTy Types;
     MapMappersArrayTy Mappers;
+    StructNonContiguousInfo NonContigInfo;
 
     /// Append arrays in \a CurInfo.
     void append(MapCombinedInfoTy &CurInfo) {
       BasePointers.append(CurInfo.BasePointers.begin(),
                           CurInfo.BasePointers.end());
       Pointers.append(CurInfo.Pointers.begin(), CurInfo.Pointers.end());
       Sizes.append(CurInfo.Sizes.begin(), CurInfo.Sizes.end());
       Types.append(CurInfo.Types.begin(), CurInfo.Types.end());
       Mappers.append(CurInfo.Mappers.begin(), CurInfo.Mappers.end());
+      NonContigInfo.Dims.append(CurInfo.NonContigInfo.Dims.begin(),
+                                 CurInfo.NonContigInfo.Dims.end());
+      NonContigInfo.Offsets.append(CurInfo.NonContigInfo.Offsets.begin(),
+                                    CurInfo.NonContigInfo.Offsets.end());
+      NonContigInfo.Counts.append(CurInfo.NonContigInfo.Counts.begin(),
+                                   CurInfo.NonContigInfo.Counts.end());
+      NonContigInfo.Strides.append(CurInfo.NonContigInfo.Strides.begin(),
+                                    CurInfo.NonContigInfo.Strides.end());
     }
   };
 
   /// Map between a struct and the its lowest & highest elements which have been
   /// mapped.
   /// [ValueDecl *] --> {LE(FieldIndex, Pointer),
   ///                    HE(FieldIndex, Pointer)}
   struct StructRangeInfoTy {
private:
 
   /// Return the corresponding bits for a given map clause modifier. Add
   /// a flag marking the map as a pointer if requested. Add a flag marking the
   /// map as the first one of a series of maps that relate to the same map
   /// expression.
   OpenMPOffloadMappingFlags getMapTypeBits(
       OpenMPMapClauseKind MapType, ArrayRef<OpenMPMapModifierKind> MapModifiers,
       ArrayRef<OpenMPMotionModifierKind> MotionModifiers, bool IsImplicit,
-      bool AddPtrFlag, bool AddIsTargetParamFlag) const {
+      bool AddPtrFlag, bool AddIsTargetParamFlag, bool IsNonContiguous) const {
     OpenMPOffloadMappingFlags Bits =
         IsImplicit ? OMP_MAP_IMPLICIT : OMP_MAP_NONE;
     switch (MapType) {
     case OMPC_MAP_alloc:
     case OMPC_MAP_release:
       // alloc and release is the default behavior in the runtime library,  i.e.
       // if we don't pass any bits alloc/release that is what the runtime is
       // going to do. Therefore, we don't need to signal anything for these two
if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_close)
         != MapModifiers.end())
       Bits |= OMP_MAP_CLOSE;
     if (llvm::find(MapModifiers, OMPC_MAP_MODIFIER_present)
         != MapModifiers.end())
       Bits |= OMP_MAP_PRESENT;
     if (llvm::find(MotionModifiers, OMPC_MOTION_MODIFIER_present)
         != MotionModifiers.end())
       Bits |= OMP_MAP_PRESENT;
+    if (IsNonContiguous)
+      Bits |= OMP_MAP_NON_CONTIG;
     return Bits;
   }
 
   /// Return true if the provided expression is a final array section. A
   /// final array section, is one whose length can't be proved to be one.
   bool isFinalArraySectionExpression(const Expr *E) const {
     const auto *OASE = dyn_cast<OMPArraySectionExpr>(E);
 
void generateInfoForComponentList(
     // in the component list which is a member expression. Useful when we have a
     // pointer or a final array section, in which case it is the previous
     // component in the list which tells us whether we have a member expression.
     // E.g. X.f[:]
     // While processing the final array section "[:]" it is "f" which tells us
     // whether we are dealing with a member of a declared struct.
     const MemberExpr *EncounteredME = nullptr;
 
+    // Track for the total number of dimension. Start from one for the dummy
+    // dimension.
+    uint64_t DimSize = 1;
+
+    bool IsNonContiguous = CombinedInfo.NonContigInfo.IsNonContiguous;
+
     for (; I != CE; ++I) {
       // If the current component is member of a struct (parent struct) mark it.
       if (!EncounteredME) {
         EncounteredME = dyn_cast<MemberExpr>(I->getAssociatedExpression());
         // If we encounter a PTR_AND_OBJ entry from now on it should be marked
         // as MEMBER_OF the parent struct.
         if (EncounteredME) {
           ShouldBeMemberOf = true;
for (; I != CE; ++I) {
       auto Next = std::next(I);
 
       // We need to generate the addresses and sizes if this is the last
       // component, if the component is a pointer or if it is an array section
       // whose length can't be proved to be one. If this is a pointer, it
       // becomes the base address for the following components.
 
       // A final array section, is one whose length can't be proved to be one.
+      // If the map item is non-contiguous then we don't treat any array section
+      // as final array section.
       bool IsFinalArraySection =
+          !IsNonContiguous &&
           isFinalArraySectionExpression(I->getAssociatedExpression());
 
       // Get information on whether the element is a pointer. Have to do a
       // special treatment for array sections given that they are built-in
       // types.
       const auto *OASE =
           dyn_cast<OMPArraySectionExpr>(I->getAssociatedExpression());
       const auto *OAShE =
           dyn_cast<OMPArrayShapingExpr>(I->getAssociatedExpression());
       const auto *UO = dyn_cast<UnaryOperator>(I->getAssociatedExpression());
       const auto *BO = dyn_cast<BinaryOperator>(I->getAssociatedExpression());
       bool IsPointer =
           OAShE ||
           (OASE && OMPArraySectionExpr::getBaseOriginalType(OASE)
                        .getCanonicalType()
                        ->isAnyPointerType()) ||
           I->getAssociatedExpression()->getType()->isAnyPointerType();
-      bool IsNonDerefPointer = IsPointer && !UO && !BO;
+      bool IsNonDerefPointer = IsPointer && !UO && !BO && !IsNonContiguous;
+
+      if (OASE)
+        ++DimSize;
 
       if (Next == CE || IsNonDerefPointer || IsFinalArraySection) {
         // If this is not the last component, we expect the pointer to be
         // associated with an array expression or member expression.
         assert((Next == CE ||
                 isa<MemberExpr>(Next->getAssociatedExpression()) ||
                 isa<ArraySubscriptExpr>(Next->getAssociatedExpression()) ||
                 isa<OMPArraySectionExpr>(Next->getAssociatedExpression()) ||
for (; I != CE; ++I) {
                   PartialStruct.HighestElem.first)>::max(),
               HB};
           PartialStruct.Base = BP;
           // Emit data for non-overlapped data.
           OpenMPOffloadMappingFlags Flags =
               OMP_MAP_MEMBER_OF |
               getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit,
                              /*AddPtrFlag=*/false,
-                             /*AddIsTargetParamFlag=*/false);
+                             /*AddIsTargetParamFlag=*/false, IsNonContiguous);
           LB = BP;
           llvm::Value *Size = nullptr;
           // Do bitcopy of all non-overlapped structure elements.
           for (OMPClauseMappableExprCommon::MappableExprComponentListRef
                    Component : OverlappedElements) {
             Address ComponentLB = Address::invalid();
             for (const OMPClauseMappableExprCommon::MappableComponent &MC :
                  Component) {
for (; I != CE; ++I) {
             }
             assert(Size && "Failed to determine structure size");
             CombinedInfo.BasePointers.push_back(BP.getPointer());
             CombinedInfo.Pointers.push_back(LB.getPointer());
             CombinedInfo.Sizes.push_back(CGF.Builder.CreateIntCast(
                 Size, CGF.Int64Ty, /*isSigned=*/true));
             CombinedInfo.Types.push_back(Flags);
             CombinedInfo.Mappers.push_back(nullptr);
+            CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize
+                                                                      : 1);
             LB = CGF.Builder.CreateConstGEP(ComponentLB, 1);
           }
           CombinedInfo.BasePointers.push_back(BP.getPointer());
           CombinedInfo.Pointers.push_back(LB.getPointer());
           Size = CGF.Builder.CreatePtrDiff(
               CGF.EmitCastToVoidPtr(
                   CGF.Builder.CreateConstGEP(HB, 1).getPointer()),
               CGF.EmitCastToVoidPtr(LB.getPointer()));
           CombinedInfo.Sizes.push_back(
               CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
           CombinedInfo.Types.push_back(Flags);
           CombinedInfo.Mappers.push_back(nullptr);
+          CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize
+                                                                    : 1);
           break;
         }
         llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression());
         if (!IsMemberPointerOrAddr) {
           CombinedInfo.BasePointers.push_back(BP.getPointer());
           CombinedInfo.Pointers.push_back(LB.getPointer());
           CombinedInfo.Sizes.push_back(
               CGF.Builder.CreateIntCast(Size, CGF.Int64Ty, /*isSigned=*/true));
+          CombinedInfo.NonContigInfo.Dims.push_back(IsNonContiguous ? DimSize
+                                                                    : 1);
 
           // If Mapper is valid, the last component inherits the mapper.
           bool HasMapper = Mapper && Next == CE;
           CombinedInfo.Mappers.push_back(HasMapper ? Mapper : nullptr);
 
           // We need to add a pointer flag for each map that comes from the
           // same expression except for the first one. We also need to signal
           // this map is the first one that relates with the current capture
           // (there is a set of entries for each capture).
-          OpenMPOffloadMappingFlags Flags =
-              getMapTypeBits(MapType, MapModifiers, MotionModifiers, IsImplicit,
-                             !IsExpressionFirstInfo || RequiresReference ||
-                                 FirstPointerInComplexData,
-                             IsCaptureFirstInfo && !RequiresReference);
+          OpenMPOffloadMappingFlags Flags = getMapTypeBits(
+              MapType, MapModifiers, MotionModifiers, IsImplicit,
+              IsCaptureFirstInfo && !RequiresReference, IsNonContiguous);
 
           if (!IsExpressionFirstInfo) {
             // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well,
             // then we reset the TO/FROM/ALWAYS/DELETE/CLOSE flags.
             if (IsPointer)
               Flags &= ~(OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_ALWAYS |
                          OMP_MAP_DELETE | OMP_MAP_CLOSE);
 
for (; I != CE; ++I) {
         if (Next != CE)
           BP = LB;
 
         IsExpressionFirstInfo = false;
         IsCaptureFirstInfo = false;
         FirstPointerInComplexData = false;
       }
     }
+
+    if (!IsNonContiguous)
+      return;
+
+    const ASTContext &Context = CGF.getContext();
+
+    // For supporting stride in array section, we need to initialize the first
+    // dimension size as 1, first offset as 0, and first count as 1
+    MapValuesArrayTy CurOffsets = {llvm::ConstantInt::get(CGF.CGM.Int64Ty, 0)};
+    MapValuesArrayTy CurCounts = {llvm::ConstantInt::get(CGF.CGM.Int64Ty, 1)};
+    MapValuesArrayTy CurStrides;
+    MapValuesArrayTy DimSizes{llvm::ConstantInt::get(CGF.CGM.Int64Ty, 1)};
+    uint64_t ElementTypeSize;
+
+    // Collect Size information for each dimension and get the element size as
+    // the first Stride. For example, for `int arr[10][10]`, the DimSizes
+    // should be [10, 10] and the first stride is 4 btyes.
+    for (const OMPClauseMappableExprCommon::MappableComponent &Component :
+         Components) {
+      const Expr *AssocExpr = Component.getAssociatedExpression();
+      const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+
+      if (!OASE)
+        continue;
+
+      QualType Ty = OMPArraySectionExpr::getBaseOriginalType(OASE->getBase());
+      auto *CAT = Context.getAsConstantArrayType(Ty);
+      auto *VAT = Context.getAsVariableArrayType(Ty);
+
+      // We need all the dimension size except for the last dimension.
+      assert((VAT || CAT || &Component == &*Components.begin()) &&
+             "Should be either ConstantArray or VariableArray if not the "
+             "first Component");
+
+      // Get element size if CurStrides is empty.
+      if (CurStrides.empty()) {
+        const Type *ElementType = nullptr;
+        if (CAT)
+          ElementType = CAT->getElementType().getTypePtr();
+        else if (VAT)
+          ElementType = VAT->getElementType().getTypePtr();
+        else
+          assert(&Component == &*Components.begin() &&
+                 "Only expect pointer (non CAT or VAT) when this is the "
+                 "first Component");
+        // If ElementType is null, then it means the base is a pointer
+        // (neither CAT nor VAT) and we'll attempt to get ElementType again
+        // for next iteration.
+        if (ElementType) {
+          // For the case that having pointer as base, we need to remove one
+          // level of indirection.
+          if (&Component != &*Components.begin())
+            ElementType = ElementType->getPointeeOrArrayElementType();
+          ElementTypeSize =
+              Context.getTypeSizeInChars(ElementType).getQuantity();
+          CurStrides.push_back(
+              llvm::ConstantInt::get(CGF.Int64Ty, ElementTypeSize));
+        }
+      }
+      // Get dimension value except for the last dimension since we don't need
+      // it.
+      if (DimSizes.size() < Components.size() - 1) {
+        if (CAT)
+          DimSizes.push_back(llvm::ConstantInt::get(
+              CGF.Int64Ty, CAT->getSize().getZExtValue()));
+        else if (VAT)
+          DimSizes.push_back(CGF.Builder.CreateIntCast(
+              CGF.EmitScalarExpr(VAT->getSizeExpr()), CGF.Int64Ty,
+              /*IsSigned=*/false));
+      }
+    }
+
+    // Skip the dummy dimension since we have already have its information.
+    auto DI = DimSizes.begin() + 1;
+    // Product of dimension.
+    llvm::Value *DimProd =
+        llvm::ConstantInt::get(CGF.CGM.Int64Ty, ElementTypeSize);
+
+    // Collect info for non-contiguous. Notice that offset, count, and stride
+    // are only meaningful for array-section, so we insert a null for anything
+    // other than array-section.
+    // Also, the size of offset, count, and stride are not the same as
+    // pointers, base_pointers, sizes, or dims. Instead, the size of offset,
+    // count, and stride are the same as the number of non-contiguous
+    // declaration in target update to/from clause.
+    for (const OMPClauseMappableExprCommon::MappableComponent &Component :
+         Components) {
+      const Expr *AssocExpr = Component.getAssociatedExpression();
+
+      if (const auto *AE = dyn_cast<ArraySubscriptExpr>(AssocExpr)) {
+        llvm::Value *Offset = CGF.Builder.CreateIntCast(
+            CGF.EmitScalarExpr(AE->getIdx()), CGF.Int64Ty,
+            /*isSigned=*/false);
+        CurOffsets.push_back(Offset);
+        CurCounts.push_back(llvm::ConstantInt::get(CGF.Int64Ty, /*V=*/1));
+        CurStrides.push_back(CurStrides.back());
+        continue;
+      }
+
+      const auto *OASE = dyn_cast<OMPArraySectionExpr>(AssocExpr);
+
+      if (!OASE)
+        continue;
+
+      // Offset
+      const Expr *OffsetExpr = OASE->getLowerBound();
+      llvm::Value *Offset = nullptr;
+      if (!OffsetExpr) {
+        // If offset is absent, then we just set it to zero.
+        Offset = llvm::ConstantInt::get(CGF.Int64Ty, 0);
+      } else {
+        Offset = CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(OffsetExpr),
+                                           CGF.Int64Ty,
+                                           /*isSigned=*/false);
+      }
+      CurOffsets.push_back(Offset);
+
+      // Count
+      const Expr *CountExpr = OASE->getLength();
+      llvm::Value *Count = nullptr;
+      if (!CountExpr) {
+        // In Clang, once a high dimension is an array section, we construct all
+        // the lower dimension as array section, however, for case like
+        // arr[0:2][2], Clang construct the inner dimension as an array section
+        // but it actually is not in an array section form according to spec.
+        if (!OASE->getColonLocFirst().isValid() &&
+            !OASE->getColonLocSecond().isValid()) {
+          Count = llvm::ConstantInt::get(CGF.Int64Ty, 1);
+        } else {
+          // OpenMP 5.0, 2.1.5 Array Sections, Description.
+          // When the length is absent it defaults to ⌈(size −
+          // lower-bound)/stride⌉, where size is the size of the array
+          // dimension.
+          const Expr *StrideExpr = OASE->getStride();
+          llvm::Value *Stride =
+              StrideExpr
+                  ? CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(StrideExpr),
+                                              CGF.Int64Ty, /*isSigned=*/false)
+                  : nullptr;
+          if (Stride)
+            Count = CGF.Builder.CreateUDiv(
+                CGF.Builder.CreateNUWSub(*DI, Offset), Stride);
+          else
+            Count = CGF.Builder.CreateNUWSub(*DI, Offset);
+        }
+      } else {
+        Count = CGF.EmitScalarExpr(CountExpr);
+      }
+      Count = CGF.Builder.CreateIntCast(Count, CGF.Int64Ty, /*isSigned=*/false);
+      CurCounts.push_back(Count);
+
+      // Stride_n' = Stride_n * (D_0 * D_1 ... * D_n-1) * Unit size
+      // Take `int arr[5][5][5]` and `arr[0:2:2][1:2:1][0:2:2]` as an example:
+      //              Offset      Count     Stride
+      //    D0          0           1         4    (int)    <- dummy dimension
+      //    D1          0           2         8    (2 * (1) * 4)
+      //    D2          1           2         20   (1 * (1 * 5) * 4)
+      //    D3          0           2         200  (2 * (1 * 5 * 4) * 4)
+      const Expr *StrideExpr = OASE->getStride();
+      llvm::Value *Stride =
+          StrideExpr
+              ? CGF.Builder.CreateIntCast(CGF.EmitScalarExpr(StrideExpr),
+                                          CGF.Int64Ty, /*isSigned=*/false)
+              : nullptr;
+      DimProd = CGF.Builder.CreateNUWMul(DimProd, *(DI - 1));
+      if (Stride)
+        CurStrides.push_back(CGF.Builder.CreateNUWMul(DimProd, Stride));
+      else
+        CurStrides.push_back(DimProd);
+      if (DI != DimSizes.end())
+        ++DI;
+    }
+
+    CombinedInfo.NonContigInfo.Offsets.push_back(CurOffsets);
+    CombinedInfo.NonContigInfo.Counts.push_back(CurCounts);
+    CombinedInfo.NonContigInfo.Strides.push_back(CurStrides);
   }
 
   /// Return the adjusted map modifiers if the declaration a capture refers to
   /// appears in a first-private clause. This is expected to be used only with
   /// directives that start with 'target'.
   MappableExprsHandler::OpenMPOffloadMappingFlags
   getMapModifiersForPrivateClauses(const CapturedStmt::Capture &Cap) const {
     assert(Cap.capturesVariable() && "Expected capture by reference only!");
for (const auto &M : Info) {
       StructRangeInfoTy PartialStruct;
 
       for (const MapInfo &L : M.second) {
         assert(!L.Components.empty() &&
                "Not expecting declaration with no component lists.");
 
         // Remember the current base pointer index.
         unsigned CurrentBasePointersIdx = CurInfo.BasePointers.size();
+        CurInfo.NonContigInfo.IsNonContiguous =
+            L.Components.back().isNonContiguous();
         generateInfoForComponentList(L.MapType, L.MapModifiers,
                                      L.MotionModifiers, L.Components, CurInfo,
                                      PartialStruct, IsFirstComponentList,
                                      L.IsImplicit, L.Mapper, L.ForDeviceAddr);
 
         // If this entry relates with a device pointer, set the relevant
         // declaration and add the 'return pointer' flag.
         if (L.ReturnDevicePointer) {
for (const auto &M : Info) {
                                      PartialStruct, IsFirstComponentList,
                                      L.IsImplicit, L.Mapper, L.ForDeviceAddr);
         IsFirstComponentList = false;
       }
 
       // If there is an entry in PartialStruct it means we have a struct with
       // individual members mapped. Emit an extra combined entry.
       if (PartialStruct.Base.isValid())
+        CurInfo.NonContigInfo.Dims.push_back(0);
         emitCombinedEntry(CombinedInfo, CurInfo.Types, PartialStruct);
 
       // We need to append the results of this capture to what we already have.
       CombinedInfo.append(CurInfo);
     }
   }
 
   /// Emit capture info for lambdas for variables captured by reference.
if (IsImplicit)
       CombinedInfo.Types.back() |= OMP_MAP_IMPLICIT;
 
     // No user-defined mapper for default mapping.
     CombinedInfo.Mappers.push_back(nullptr);
   }
 };
 } // anonymous namespace
 
+static void emitNonContiguousDescriptor(
+    CodeGenFunction &CGF, MappableExprsHandler::MapCombinedInfoTy &CombinedInfo,
+    CGOpenMPRuntime::TargetDataInfo &Info) {
+  CodeGenModule &CGM = CGF.CGM;
+  MappableExprsHandler::MapCombinedInfoTy::StructNonContiguousInfo
+      &NonContigInfo = CombinedInfo.NonContigInfo;
+
+  // Build an array of struct descriptor_dim and then assign it to
+  // offload_args.
+  //
+  // struct descriptor_dim {
+  //  uint64_t offset;
+  //  uint64_t count;
+  //  uint64_t stride
+  // };
+  ASTContext &C = CGF.getContext();
+  QualType Int64Ty = C.getIntTypeForBitwidth(/*DestWidth=*/64, /*Signed=*/0);
+  RecordDecl *RD;
+  RD = C.buildImplicitRecord("descriptor_dim");
+  RD->startDefinition();
+  addFieldToRecordDecl(C, RD, Int64Ty);
+  addFieldToRecordDecl(C, RD, Int64Ty);
+  addFieldToRecordDecl(C, RD, Int64Ty);
+  RD->completeDefinition();
+  QualType DimTy = C.getRecordType(RD);
+
+  enum { OffsetFD = 0, CountFD, StrideFD };
+  // We need two index variable here since the size of "Dims" is the same as the
+  // size of Components, however, the size of offset, count, and stride is equal
+  // to the size of base declaration that is non-contiguous.
+  for (unsigned I = 0, L = 0, E = NonContigInfo.Dims.size(); I < E; ++I) {
+    // Skip emitting ir if dimension size is 1 since it cannot be
+    // non-contiguous.
+    if (NonContigInfo.Dims[I] == 1)
+      continue;
+    llvm::APInt Size(/*numBits=*/32, NonContigInfo.Dims[I]);
+    QualType ArrayTy =
+        C.getConstantArrayType(DimTy, Size, nullptr, ArrayType::Normal, 0);
+    Address DimsAddr = CGF.CreateMemTemp(ArrayTy, "dims");
+    for (unsigned II = 0, EE = NonContigInfo.Dims[I]; II < EE; ++II) {
+      unsigned RevIdx = EE - II - 1;
+      LValue DimsLVal = CGF.MakeAddrLValue(
+          CGF.Builder.CreateConstArrayGEP(DimsAddr, II), DimTy);
+      // Offset
+      LValue OffsetLVal = CGF.EmitLValueForField(
+          DimsLVal, *std::next(RD->field_begin(), OffsetFD));
+      CGF.EmitStoreOfScalar(NonContigInfo.Offsets[L][RevIdx], OffsetLVal);
+      // Count
+      LValue CountLVal = CGF.EmitLValueForField(
+          DimsLVal, *std::next(RD->field_begin(), CountFD));
+      CGF.EmitStoreOfScalar(NonContigInfo.Counts[L][RevIdx], CountLVal);
+      // Stride
+      LValue StrideLVal = CGF.EmitLValueForField(
+          DimsLVal, *std::next(RD->field_begin(), StrideFD));
+      CGF.EmitStoreOfScalar(NonContigInfo.Strides[L][RevIdx], StrideLVal);
+    }
+    // args[I] = &dims
+    Address DAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
+        DimsAddr, CGM.Int8PtrTy);
+    llvm::Value *P = CGF.Builder.CreateConstInBoundsGEP2_32(
+        llvm::ArrayType::get(CGM.VoidPtrTy, Info.NumberOfPtrs),
+        Info.PointersArray, 0, I);
+    Address PAddr(P, CGF.getPointerAlign());
+    CGF.Builder.CreateStore(DAddr.getPointer(), PAddr);
+    ++L;
+  }
+}
+
 /// Emit the arrays used to pass the captures and map information to the
 /// offloading runtime library. If there is no map or capture information,
 /// return nullptr by reference.
 static void
 emitOffloadingArrays(CodeGenFunction &CGF,
                      MappableExprsHandler::MapCombinedInfoTy &CombinedInfo,
-                     CGOpenMPRuntime::TargetDataInfo &Info) {
+                     CGOpenMPRuntime::TargetDataInfo &Info,
+                     bool IsNonContiguous = false) {
   CodeGenModule &CGM = CGF.CGM;
   ASTContext &Ctx = CGF.getContext();
 
   // Reset the array information.
   Info.clearArrayInfo();
   Info.NumberOfPtrs = CombinedInfo.BasePointers.size();
 
   if (Info.NumberOfPtrs) {
if (hasRuntimeEvaluationCaptureSize) {
           Int64Ty, PointerNumAP, nullptr, ArrayType::Normal,
           /*IndexTypeQuals=*/0);
       Info.SizesArray =
           CGF.CreateMemTemp(SizeArrayType, ".offload_sizes").getPointer();
     } else {
       // We expect all the sizes to be constant, so we collect them to create
       // a constant array.
       SmallVector<llvm::Constant *, 16> ConstSizes;
-      for (llvm::Value *S : CombinedInfo.Sizes)
-        ConstSizes.push_back(cast<llvm::Constant>(S));
+      for (unsigned I = 0, E = CombinedInfo.Sizes.size(); I < E; ++I) {
+        if (IsNonContiguous &&
+            (CombinedInfo.Types[I] & MappableExprsHandler::OMP_MAP_NON_CONTIG)) {
+          ConstSizes.push_back(llvm::ConstantInt::get(
+              CGF.Int64Ty, CombinedInfo.NonContigInfo.Dims[I]));
+        } else {
+          ConstSizes.push_back(cast<llvm::Constant>(CombinedInfo.Sizes[I]));
+        }
+      }
 
       auto *SizesArrayInit = llvm::ConstantArray::get(
           llvm::ArrayType::get(CGM.Int64Ty, ConstSizes.size()), ConstSizes);
       std::string Name = CGM.getOpenMPRuntime().getName({"offload_sizes"});
       auto *SizesArrayGbl = new llvm::GlobalVariable(
           CGM.getModule(), SizesArrayInit->getType(),
           /*isConstant=*/true, llvm::GlobalValue::PrivateLinkage,
           SizesArrayInit, Name);
for (unsigned I = 0; I < Info.NumberOfPtrs; ++I) {
             cast<OMPDeclareMapperDecl>(CombinedInfo.Mappers[I]));
         MFunc = CGF.Builder.CreatePointerCast(MFunc, CGM.VoidPtrTy);
         Info.HasMapper = true;
       }
       Address MAddr = CGF.Builder.CreateConstArrayGEP(MappersArray, I);
       CGF.Builder.CreateStore(MFunc, MAddr);
     }
   }
+
+  if (!IsNonContiguous || CombinedInfo.NonContigInfo.Offsets.empty() ||
+      Info.NumberOfPtrs == 0)
+    return;
+
+  emitNonContiguousDescriptor(CGF, CombinedInfo, Info);
 }
 
 namespace {
 /// Additional arguments for emitOffloadingArraysArgument function.
 struct ArgumentsOptions {
   bool ForEndCall = false;
   ArgumentsOptions() = default;
   ArgumentsOptions(bool ForEndCall) : ForEndCall(ForEndCall) {}
auto &&BeginThenGen = [this, &D, Device, &Info,
     // Fill up the arrays with all the mapped variables.
     MappableExprsHandler::MapCombinedInfoTy CombinedInfo;
 
     // Get map clause information.
     MappableExprsHandler MEHandler(D, CGF);
     MEHandler.generateAllInfo(CombinedInfo);
 
     // Fill up the arrays and create the arguments.
-    emitOffloadingArrays(CGF, CombinedInfo, Info);
+    emitOffloadingArrays(CGF, CombinedInfo, Info, /*IsNonContiguous=*/true);
 
     llvm::Value *BasePointersArrayArg = nullptr;
     llvm::Value *PointersArrayArg = nullptr;
     llvm::Value *SizesArrayArg = nullptr;
     llvm::Value *MapTypesArrayArg = nullptr;
     llvm::Value *MappersArrayArg = nullptr;
     emitOffloadingArraysArgument(CGF, BasePointersArrayArg, PointersArrayArg,
                                  SizesArrayArg, MapTypesArrayArg,
auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray](
     MappableExprsHandler::MapCombinedInfoTy CombinedInfo;
 
     // Get map clause information.
     MappableExprsHandler MEHandler(D, CGF);
     MEHandler.generateAllInfo(CombinedInfo);
 
     TargetDataInfo Info;
     // Fill up the arrays and create the arguments.
-    emitOffloadingArrays(CGF, CombinedInfo, Info);
+    emitOffloadingArrays(CGF, CombinedInfo, Info, /*IsNonContiguous=*/true);
     bool RequiresOuterTask = D.hasClausesOfKind<OMPDependClause>() ||
                              D.hasClausesOfKind<OMPNowaitClause>();
     emitOffloadingArraysArgument(
         CGF, Info.BasePointersArray, Info.PointersArray, Info.SizesArray,
         Info.MapTypesArray, Info.MappersArray, Info, {/*ForEndTask=*/false});
     InputInfo.NumberOfTargetItems = Info.NumberOfPtrs;
     InputInfo.BasePointersArray =
         Address(Info.BasePointersArray, CGM.getPointerAlign());

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'.
   DSA_firstprivate = 1 << 2, /// Default data sharing attribute 'firstprivate'.
if (auto *TE = dyn_cast<CXXThisExpr>(E->getBase()->IgnoreParenCasts())) {
         if (DVar.CKind != OMPC_unknown)
           ImplicitFirstprivate.push_back(E);
       }
       return;
     }
     if (isOpenMPTargetExecutionDirective(DKind)) {
       OMPClauseMappableExprCommon::MappableExprComponentList CurComponents;
       if (!checkMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map,
+                                        Stack->getCurrentDirective(),
                                         /*NoDiagnose=*/true))
         return;
       const auto *VD = cast<ValueDecl>(
           CurComponents.back().getAssociatedDeclaration()->getCanonicalDecl());
       if (!Stack->checkMappableExprComponentListsForDecl(
               VD, /*CurrentRegionOnly=*/true,
               [&CurComponents](
                   OMPClauseMappableExprCommon::MappableExprComponentListRef

 // but these would be valid:
 //   r.ArrS[3].Arr[6:7]
 //
 //   r.ArrS[3].x
 namespace {
 class MapBaseChecker final : public StmtVisitor<MapBaseChecker, bool> {
   Sema &SemaRef;
   OpenMPClauseKind CKind = OMPC_unknown;
+  OpenMPDirectiveKind DKind = OMPD_unknown;
   OMPClauseMappableExprCommon::MappableExprComponentList &Components;
+  bool IsNonContiguous = false;
   bool NoDiagnose = false;
   const Expr *RelevantExpr = nullptr;
   bool AllowUnitySizeArraySection = true;
   bool AllowWholeSizeArraySection = true;
+  bool AllowAnotherPtr = true;
   SourceLocation ELoc;
   SourceRange ERange;
 
   void emitErrorMsg() {
     // If nothing else worked, this is not a valid map clause expression.
     if (SemaRef.getLangOpts().OpenMP < 50) {
       SemaRef.Diag(ELoc,
                    diag::err_omp_expected_named_var_member_or_array_expression)
           << ERange;
     } else {
       SemaRef.Diag(ELoc, diag::err_omp_non_lvalue_in_map_or_motion_clauses)
           << getOpenMPClauseName(CKind) << ERange;
     }
   }
 
 public:
   bool VisitDeclRefExpr(DeclRefExpr *DRE) {
     if (!isa<VarDecl>(DRE->getDecl())) {
       emitErrorMsg();
       return false;
     }
     assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
     RelevantExpr = DRE;
     // Record the component.
-    Components.emplace_back(DRE, DRE->getDecl());
+    Components.emplace_back(DRE, DRE->getDecl(), IsNonContiguous);
     return true;
   }
 
   bool VisitMemberExpr(MemberExpr *ME) {
     Expr *E = ME;
     Expr *BaseE = ME->getBase()->IgnoreParenCasts();
 
     if (isa<CXXThisExpr>(BaseE)) {
bool VisitMemberExpr(MemberExpr *ME) {
     // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.7]
     //  If a list item is an element of a structure, only the rightmost symbol
     //  of the variable reference can be an array section.
     //
     AllowUnitySizeArraySection = false;
     AllowWholeSizeArraySection = false;
 
     // Record the component.
-    Components.emplace_back(ME, FD);
+    Components.emplace_back(ME, FD, IsNonContiguous);
     return RelevantExpr || Visit(E);
   }
 
   bool VisitArraySubscriptExpr(ArraySubscriptExpr *AE) {
     Expr *E = AE->getBase()->IgnoreParenImpCasts();
 
     if (!E->getType()->isAnyPointerType() && !E->getType()->isArrayType()) {
       if (!NoDiagnose) {
if (const auto *TE = dyn_cast<CXXThisExpr>(E->IgnoreParenCasts())) {
         SemaRef.Diag(AE->getIdx()->getExprLoc(),
                      diag::note_omp_invalid_subscript_on_this_ptr_map);
       }
       assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
       RelevantExpr = TE;
     }
 
     // Record the component - we don't have any declaration associated.
-    Components.emplace_back(AE, nullptr);
+    Components.emplace_back(AE, nullptr, IsNonContiguous);
 
     return RelevantExpr || Visit(E);
   }
 
   bool VisitOMPArraySectionExpr(OMPArraySectionExpr *OASE) {
     assert(!NoDiagnose && "Array sections cannot be implicitly mapped.");
     Expr *E = OASE->getBase()->IgnoreParenImpCasts();
     QualType 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 (DKind == OMPD_target_update &&
+               SemaRef.getLangOpts().OpenMP >= 50) {
+      if (IsPointer && !AllowAnotherPtr)
+        SemaRef.Diag(ELoc, diag::err_omp_section_length_undefined)
+            << /*array of unknown bound */ 1;
+      else
+        IsNonContiguous = true;
     } else if (AllowUnitySizeArraySection && NotUnity) {
       // A unity or whole array section is not allowed and that is not
       // compatible with the properties of the current array section.
       SemaRef.Diag(
         ELoc, diag::err_array_section_does_not_specify_contiguous_storage)
         << OASE->getSourceRange();
       return false;
     }
 
+    if (IsPointer)
+      AllowAnotherPtr = false;
+
     if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
       Expr::EvalResult ResultR;
       Expr::EvalResult ResultL;
       if (!OASE->getLength()->isValueDependent() &&
           OASE->getLength()->EvaluateAsInt(ResultR, SemaRef.getASTContext()) &&
           !ResultR.Val.getInt().isOneValue()) {
         SemaRef.Diag(OASE->getLength()->getExprLoc(),
                      diag::err_omp_invalid_map_this_expr);
if (const auto *TE = dyn_cast<CXXThisExpr>(E)) {
         SemaRef.Diag(OASE->getLowerBound()->getExprLoc(),
                      diag::note_omp_invalid_lower_bound_on_this_ptr_mapping);
       }
       assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
       RelevantExpr = TE;
     }
 
     // Record the component - we don't have any declaration associated.
-    Components.emplace_back(OASE, nullptr);
+    Components.emplace_back(OASE, nullptr, /*IsNonContiguous=*/false);
     return RelevantExpr || Visit(E);
   }
   bool VisitOMPArrayShapingExpr(OMPArrayShapingExpr *E) {
     Expr *Base = E->getBase();
 
     // Record the component - we don't have any declaration associated.
-    Components.emplace_back(E, nullptr);
+    Components.emplace_back(E, nullptr, IsNonContiguous);
 
     return Visit(Base->IgnoreParenImpCasts());
   }
 
   bool VisitUnaryOperator(UnaryOperator *UO) {
     if (SemaRef.getLangOpts().OpenMP < 50 || !UO->isLValue() ||
         UO->getOpcode() != UO_Deref) {
       emitErrorMsg();
       return false;
     }
     if (!RelevantExpr) {
       // Record the component if haven't found base decl.
-      Components.emplace_back(UO, nullptr);
+      Components.emplace_back(UO, nullptr, /*IsNonContiguous=*/false);
     }
     return RelevantExpr || Visit(UO->getSubExpr()->IgnoreParenImpCasts());
   }
   bool VisitBinaryOperator(BinaryOperator *BO) {
     if (SemaRef.getLangOpts().OpenMP < 50 || !BO->getType()->isPointerType()) {
       emitErrorMsg();
       return false;
     }
 
     // Pointer arithmetic is the only thing we expect to happen here so after we
     // make sure the binary operator is a pointer type, the we only thing need
     // to to is to visit the subtree that has the same type as root (so that we
     // know the other subtree is just an offset)
     Expr *LE = BO->getLHS()->IgnoreParenImpCasts();
     Expr *RE = BO->getRHS()->IgnoreParenImpCasts();
-    Components.emplace_back(BO, nullptr);
+    Components.emplace_back(BO, nullptr, false);
     assert((LE->getType().getTypePtr() == BO->getType().getTypePtr() ||
             RE->getType().getTypePtr() == BO->getType().getTypePtr()) &&
            "Either LHS or RHS have base decl inside");
     if (BO->getType().getTypePtr() == LE->getType().getTypePtr())
       return RelevantExpr || Visit(LE);
     return RelevantExpr || Visit(RE);
   }
   bool VisitCXXThisExpr(CXXThisExpr *CTE) {
     assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
     RelevantExpr = CTE;
-    Components.emplace_back(CTE, nullptr);
+    Components.emplace_back(CTE, nullptr, IsNonContiguous);
     return true;
   }
   bool VisitCXXOperatorCallExpr(CXXOperatorCallExpr *COCE) {
     assert(!RelevantExpr && "RelevantExpr is expected to be nullptr");
-    Components.emplace_back(COCE, nullptr);
+    Components.emplace_back(COCE, nullptr, IsNonContiguous);
     return true;
   }
   bool VisitStmt(Stmt *) {
     emitErrorMsg();
     return false;
   }
   const Expr *getFoundBase() const {
     return RelevantExpr;
   }
   explicit MapBaseChecker(
-      Sema &SemaRef, OpenMPClauseKind CKind,
+      Sema &SemaRef, OpenMPClauseKind CKind, OpenMPDirectiveKind DKind,
       OMPClauseMappableExprCommon::MappableExprComponentList &Components,
       bool NoDiagnose, SourceLocation &ELoc, SourceRange &ERange)
-      : SemaRef(SemaRef), CKind(CKind), Components(Components),
+      : SemaRef(SemaRef), CKind(CKind), DKind(DKind), Components(Components),
         NoDiagnose(NoDiagnose), ELoc(ELoc), ERange(ERange) {}
 };
 } // namespace
 
 /// Return the expression of the base of the mappable expression or null if it
 /// cannot be determined and do all the necessary checks to see if the expression
 /// is valid as a standalone mappable expression. In the process, record all the
 /// components of the expression.
 static const Expr *checkMapClauseExpressionBase(
     Sema &SemaRef, Expr *E,
     OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents,
-    OpenMPClauseKind CKind, bool NoDiagnose) {
+    OpenMPClauseKind CKind, OpenMPDirectiveKind DKind, bool NoDiagnose) {
   SourceLocation ELoc = E->getExprLoc();
   SourceRange ERange = E->getSourceRange();
-  MapBaseChecker Checker(SemaRef, CKind, CurComponents, NoDiagnose, ELoc,
+  MapBaseChecker Checker(SemaRef, CKind, DKind, CurComponents, NoDiagnose, ELoc,
                          ERange);
-  if (Checker.Visit(E->IgnoreParens()))
+  if (Checker.Visit(E->IgnoreParens())) {
+    // Check if the highest dimension array section has length specified
+    if (SemaRef.getLangOpts().OpenMP >= 50 && !CurComponents.empty() &&
+        (CKind == OMPC_to || CKind == OMPC_from)) {
+      auto CI = CurComponents.rbegin();
+      auto CE = CurComponents.rend();
+      for (; CI != CE; ++CI) {
+        const auto *OASE =
+            dyn_cast<OMPArraySectionExpr>(CI->getAssociatedExpression());
+        if (!OASE)
+          continue;
+        if (OASE && OASE->getLength())
+          break;
+        SemaRef.Diag(ELoc, diag::err_array_section_does_not_specify_length)
+            << ERange;
+      }
+    }
     return Checker.getFoundBase();
+  }
   return nullptr;
 }
 
 // Return true if expression E associated with value VD has conflicts with other
 // map information.
 static bool checkMapConflicts(
     Sema &SemaRef, DSAStackTy *DSAS, const ValueDecl *VD, const Expr *E,
     bool CurrentRegionOnly,
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
for (Expr *RefExpr : VarList) {
     // is correctly captured. A variable that shows up in a use_device_ptr has
     // similar properties of a first private variable.
     DSAStack->addDSA(D, RefExpr->IgnoreParens(), OMPC_firstprivate, Ref);
 
     // Create a mappable component for the list item. List items in this clause
     // only need a component.
     MVLI.VarBaseDeclarations.push_back(D);
     MVLI.VarComponents.resize(MVLI.VarComponents.size() + 1);
-    MVLI.VarComponents.back().push_back(
-        OMPClauseMappableExprCommon::MappableComponent(SimpleRefExpr, D));
+    MVLI.VarComponents.back().emplace_back(SimpleRefExpr, D,
+                                           /*IsNonContiguous=*/false);
   }
 
   if (MVLI.ProcessedVarList.empty())
     return nullptr;
 
   return OMPUseDevicePtrClause::Create(
       Context, Locs, MVLI.ProcessedVarList, PrivateCopies, Inits,
       MVLI.VarBaseDeclarations, MVLI.VarComponents);
for (Expr *RefExpr : VarList) {
     // Create a mappable component for the list item. List items in this clause
     // only need a component.
     MVLI.VarBaseDeclarations.push_back(D);
     MVLI.VarComponents.emplace_back();
     Expr *Component = SimpleRefExpr;
     if (VD && (isa<OMPArraySectionExpr>(RefExpr->IgnoreParenImpCasts()) ||
                isa<ArraySubscriptExpr>(RefExpr->IgnoreParenImpCasts())))
       Component = DefaultFunctionArrayLvalueConversion(SimpleRefExpr).get();
-    MVLI.VarComponents.back().push_back(
-        OMPClauseMappableExprCommon::MappableComponent(Component, D));
+    MVLI.VarComponents.back().emplace_back(Component, D,
+                                           /*IsNonContiguous=*/false);
   }
 
   if (MVLI.ProcessedVarList.empty())
     return nullptr;
 
   return OMPUseDeviceAddrClause::Create(Context, Locs, MVLI.ProcessedVarList,
                                         MVLI.VarBaseDeclarations,
                                         MVLI.VarComponents);
if (DSAStack->checkMappableExprComponentListsForDecl(
       Diag(ELoc, diag::err_omp_map_shared_storage) << RefExpr->getSourceRange();
       Diag(ConflictExpr->getExprLoc(), diag::note_used_here)
           << ConflictExpr->getSourceRange();
       continue;
     }
 
     // Store the components in the stack so that they can be used to check
     // against other clauses later on.
-    OMPClauseMappableExprCommon::MappableComponent MC(SimpleRefExpr, D);
+    OMPClauseMappableExprCommon::MappableComponent MC(
+        SimpleRefExpr, D, /*IsNonContiguous=*/false);
     DSAStack->addMappableExpressionComponents(
         D, MC, /*WhereFoundClauseKind=*/OMPC_is_device_ptr);
 
     // Record the expression we've just processed.
     MVLI.ProcessedVarList.push_back(SimpleRefExpr);
 
     // Create a mappable component for the list item. List items in this clause
     // only need a component. We use a null declaration to signal fields in

clang/lib/Serialization/ASTReader.cpp

void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) {
   ListSizes.reserve(TotalLists);
   for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
-    Expr *AssociatedExpr = Record.readExpr();
+    Expr *AssociatedExprPr = Record.readExpr();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(AssociatedExprPr, AssociatedDecl,
+                            /*IsNonContiguous=*/false);
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPAllocateClause(OMPAllocateClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   C->setColonLoc(Record.readSourceLocation());
   C->setAllocator(Record.readSubExpr());
void OMPClauseReader::VisitOMPToClause(OMPToClause *C) {
   ListSizes.reserve(TotalLists);
   for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
-    Expr *AssociatedExpr = Record.readSubExpr();
+    Expr *AssociatedExprPr = Record.readSubExpr();
+    bool IsNonContiguous = Record.readBool();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(AssociatedExprPr, AssociatedDecl, IsNonContiguous);
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) {
     C->setMotionModifier(
void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) {
   ListSizes.reserve(TotalLists);
   for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
-    Expr *AssociatedExpr = Record.readSubExpr();
+    Expr *AssociatedExprPr = Record.readSubExpr();
+    bool IsNonContiguous = Record.readBool();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(AssociatedExprPr, AssociatedDecl, IsNonContiguous);
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   auto NumVars = C->varlist_size();
   auto UniqueDecls = C->getUniqueDeclarationsNum();
void OMPClauseReader::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   ListSizes.reserve(TotalLists);
   for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
-    Expr *AssociatedExpr = Record.readSubExpr();
+    auto *AssociatedExprPr = Record.readSubExpr();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(AssociatedExprPr, AssociatedDecl,
+                            /*IsNonContiguous=*/false);
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPUseDeviceAddrClause(OMPUseDeviceAddrClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   auto NumVars = C->varlist_size();
   auto UniqueDecls = C->getUniqueDeclarationsNum();
for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
     Expr *AssociatedExpr = Record.readSubExpr();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(AssociatedExpr, AssociatedDecl,
+                            /*IsNonContiguous*/ false);
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPIsDevicePtrClause(OMPIsDevicePtrClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   auto NumVars = C->varlist_size();
   auto UniqueDecls = C->getUniqueDeclarationsNum();
for (unsigned i = 0; i < TotalLists; ++i)
     ListSizes.push_back(Record.readInt());
   C->setComponentListSizes(ListSizes);
 
   SmallVector<OMPClauseMappableExprCommon::MappableComponent, 32> Components;
   Components.reserve(TotalComponents);
   for (unsigned i = 0; i < TotalComponents; ++i) {
     Expr *AssociatedExpr = Record.readSubExpr();
     auto *AssociatedDecl = Record.readDeclAs<ValueDecl>();
-    Components.push_back(OMPClauseMappableExprCommon::MappableComponent(
-        AssociatedExpr, AssociatedDecl));
+    Components.emplace_back(AssociatedExpr, AssociatedDecl,
+                            /*IsNonContiguous=*/false);
   }
   C->setComponents(Components, ListSizes);
 }
 
 void OMPClauseReader::VisitOMPNontemporalClause(OMPNontemporalClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
   unsigned NumVars = C->varlist_size();
   SmallVector<Expr *, 16> Vars;

clang/lib/Serialization/ASTWriter.cpp

void OMPClauseWriter::VisitOMPToClause(OMPToClause *C) {
   for (auto *D : C->all_decls())
     Record.AddDeclRef(D);
   for (auto N : C->all_num_lists())
     Record.push_back(N);
   for (auto N : C->all_lists_sizes())
     Record.push_back(N);
   for (auto &M : C->all_components()) {
     Record.AddStmt(M.getAssociatedExpression());
+    Record.writeBool(M.isNonContiguous());
     Record.AddDeclRef(M.getAssociatedDeclaration());
   }
 }
 
 void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) {
   Record.push_back(C->varlist_size());
   Record.push_back(C->getUniqueDeclarationsNum());
   Record.push_back(C->getTotalComponentListNum());
void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) {
   for (auto *D : C->all_decls())
     Record.AddDeclRef(D);
   for (auto N : C->all_num_lists())
     Record.push_back(N);
   for (auto N : C->all_lists_sizes())
     Record.push_back(N);
   for (auto &M : C->all_components()) {
     Record.AddStmt(M.getAssociatedExpression());
+    Record.writeBool(M.isNonContiguous());
     Record.AddDeclRef(M.getAssociatedDeclaration());
   }
 }
 
 void OMPClauseWriter::VisitOMPUseDevicePtrClause(OMPUseDevicePtrClause *C) {
   Record.push_back(C->varlist_size());
   Record.push_back(C->getUniqueDeclarationsNum());
   Record.push_back(C->getTotalComponentListNum());

clang/test/OpenMP/target_update_ast_print.cpp

 T foo(T targ, U uarg) {
   static T a, *p;
   U b;
   int l;
 #pragma omp target update to(([a][targ])p, a) if(l>5) device(l) nowait depend(inout:l)
 
 #pragma omp target update from(b, ([a][targ])p) 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 from(marr[:2][0:2][0:2:1])
+
+#pragma omp target update to(marr[:l][: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])
+
   int arr[100][100];
 #pragma omp target update to(arr[2][0:1:2])
 
 #pragma omp target update from(arr[2][0:1:2])
 
 #ifdef OMP51
 #pragma omp target update to(present: arr[2][0:1:2])
 

 // CHECK-NEXT: int l;
 // CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l)
 // CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l)
 // CHECK:      static char a, *p;
 // CHECK-NEXT: float b;
 // CHECK-NEXT: int l;
 // CHECK-NEXT: #pragma omp target update to(([a][targ])p,a) if(l > 5) device(l) nowait depend(inout : l)
 // CHECK-NEXT: #pragma omp target update from(b,([a][targ])p) if(l < 5) device(l - 1) nowait depend(inout : l)
+// CHECK: 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 from(marr[:2][0:2][0:2:1])
+// CHECK-NEXT: #pragma omp target update to(marr[:l][: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])
 // CHECK:      int arr[100][100];
 // CHECK-NEXT: #pragma omp target update to(arr[2][0:1:2])
 // CHECK-NEXT: #pragma omp target update from(arr[2][0:1:2])
 // OMP5-NEXT: #pragma omp target update to(present: arr[2][0:1:2])
 // OMP5-NEXT: #pragma omp target update from(present: arr[2][0:1:2], a)
 
 int main(int argc, char **argv) {
   static int a;

 // CHECK-NEXT: #pragma omp target update from(argv[2][0:1:2])
 #ifdef OMP51
 #pragma omp target update to(present: argv[2][0:1:2])
 // OMP5-NEXT: #pragma omp target update to(present: arr[2][0:1:2])
 #pragma omp target update from(argv[2][0:1:2], a)
 // OMP5-NEXT: #pragma omp target update from(present: arr[2][0:1:2], a)
 #endif
 
+float marr[10][10][10];
+// CHECK: 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[:n][:n][n:])
+// CHECK: #pragma omp target update to(marr[:n][: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])
+#pragma omp target update to(marr[:2:][0:2][0:2:1])
+// CHECK-NEXT: #pragma omp target update to(marr[:2:][0:2][0:2:1])
+#pragma omp target update from(marr[:2:][0:2][0:2:1])
+// CHECK-NEXT: #pragma omp target update from(marr[:2:][0:2][0:2:1])
+#pragma omp target update to(marr[:2:][:2:][0:2:1])
+// CHECK-NEXT: #pragma omp target update to(marr[:2:][:2:][0:2:1])
+#pragma omp target update from(marr[:2:][:2:][0:2:1])
+// CHECK-NEXT: #pragma omp target update from(marr[:2:][:2:][0:2:1])
 
   return foo(argc, f) + foo(argv[0][0], f) + a;
 }
 
 #endif

clang/test/OpenMP/target_update_codegen.cpp

void check_present(int arg) {
   // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
   // CK19-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
   // CK19-DAG: store float* [[VAL0]], float** [[CP0]]
   // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
   #pragma omp target update from(present: lb)
   ;
 }
 #endif
+
+///==========================================================================///
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK20 --check-prefix CK20-64
+// RUN: %clang_cc1 -DCK20 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK20 --check-prefix CK20-64
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK20 --check-prefix CK20-32
+// RUN: %clang_cc1 -DCK20 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK20 --check-prefix CK20-32
+
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK20 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK20 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK20
+
+struct ST {
+  int a;
+  double *b;
+};
+
+// CK20: [[STRUCT_ST:%.+]] = type { i32, double* }
+// CK20: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK20: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 3]
+// CK20: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+// CK20-LABEL: _Z3foo
+void foo(int arg) {
+  ST arr[3][4];
+  // CK20: [[DIMS:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]],
+  // CK20: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [[STRUCT_ST]]]], [3 x [4 x [[STRUCT_ST]]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+  // CK20: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [[STRUCT_ST]]], [4 x [[STRUCT_ST]]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+  // CK20: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [[STRUCT_ST]], [[STRUCT_ST]]* [[ARRAY_DECAY]], {{.+}}
+  // CK20: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+  // CK20: [[BPC:%.+]] = bitcast i8** [[BP0]] to [3 x [4 x [[STRUCT_ST]]]]**
+  // CK20: store [3 x [4 x [[STRUCT_ST]]]]* [[ARR]], [3 x [4 x [[STRUCT_ST]]]]** [[BPC]],
+  // CK20: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0
+  // CK20: [[PC:%.+]] = bitcast i8** [[P0]] to [[STRUCT_ST]]**
+  // CK20: store [[STRUCT_ST]]* [[ARRAY_IDX_1]], [[STRUCT_ST]]** [[PC]],
+  // CK20: [[DIM_1:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+  // CK20: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+  // CK20: store i64 0, i64* [[OFFSET]],
+  // CK20: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+  // CK20: store i64 2, i64* [[COUNT]],
+  // CK20: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+  // CK20: store i64 {{32|64}}, i64* [[STRIDE]],
+  // CK20: [[DIM_2:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+  // CK20: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+  // CK20: store i64 1, i64* [[OFFSET_2]],
+  // CK20: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+  // CK20: store i64 4, i64* [[COUNT_2]],
+  // CK20: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+  // CK20: store i64 {{8|16}}, i64* [[STRIDE_2]],
+  // CK20: [[DIM_3:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+  // CK20: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+  // CK20: store i64 0, i64* [[OFFSET_3]],
+  // CK20: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+  // CK20: store i64 1, i64* [[COUNT_3]],
+  // CK20: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+  // CK20: store i64 {{8|16}}, i64* [[STRIDE_3]],
+  // CK20-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+  // CK20-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK20-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK20-DAG: [[PC0:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+  // CK20-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+  // CK20-DAG: store i8* [[PC0]], i8** [[PTRS]],
+
+#pragma omp target update to(arr[0:2][1:4])
+  { ++arg; }
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK21 --check-prefix CK21-64
+// RUN: %clang_cc1 -DCK21 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK21 --check-prefix CK21-64
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK21 --check-prefix CK21-32
+// RUN: %clang_cc1 -DCK21 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK21 --check-prefix CK21-32
+
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK21 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK21 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK21
+
+// CK21: [[STRUCT_ST:%.+]] = type { [10 x [10 x [10 x double*]]] }
+// CK21: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK21: [[MTYPE:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 299067162755073]
+
+struct ST {
+  double *dptr[10][10][10];
+
+  // CK21: _ZN2ST3fooEv
+  void foo() {
+    // CK21: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+    // CK21: [[ARRAY_IDX:%.+]] = getelementptr inbounds [10 x [10 x [10 x double*]]], [10 x [10 x [10 x double*]]]* [[DPTR:%.+]], {{.+}} 0, {{.+}} 0
+    // CK21: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [10 x [10 x double*]], [10 x [10 x double*]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+    // CK21: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [10 x double*], [10 x double*]* [[ARRAY_DECAY]], {{.+}} 1
+    // CK21: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [10 x double*], [10 x double*]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+    // CK21: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds {{.+}}, {{.+}}* [[ARRAY_DECAY_2]], {{.+}} 0
+    // CK21: [[BP0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+    // CK21: [[P0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0
+    // CK21: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+    // CK21: [[OFFSET_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+    // CK21: store i64 0, i64* [[OFFSET_1]],
+    // CK21: [[COUNT_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+    // CK21: store i64 2, i64* [[COUNT_1]],
+    // CK21: [[STRIDE_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+    // CK21: store i64 {{400|800}}, i64* [[STRIDE_1]],
+    // CK21: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+    // CK21: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+    // CK21: store i64 1, i64* [[OFFSET_2]],
+    // CK21: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+    // CK21: store i64 3, i64* [[COUNT_2]],
+    // CK21: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+    // CK21: store i64 {{40|80}}, i64* [[STRIDE_2]],
+    // CK21: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+    // CK21: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+    // CK21: store i64 0, i64* [[OFFSET_3]],
+    // CK21: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+    // CK21: store i64 4, i64* [[COUNT_3]],
+    // CK21: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+    // CK21: store i64 {{4|8}}, i64* [[STRIDE_3]],
+    // CK21: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+    // CK21: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+    // CK21: store i64 0, i64* [[OFFSET_4]],
+    // CK21: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+    // CK21: store i64 1, i64* [[COUNT_4]],
+    // CK21: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+    // CK21: store i64 {{4|8}}, i64* [[STRIDE_4]],
+    // CK21-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}}* [[GEPSZ:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE]]{{.+}})
+    // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+    // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+    // CK21-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+    // CK21-DAG: [[PTRS:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i32 0, i32 0
+    // CK21-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(dptr[0:2][1:3][0:4])
+  }
+};
+
+void bar() {
+  ST st;
+  st.foo();
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK22 --check-prefix CK22-64
+// RUN: %clang_cc1 -DCK22 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK22 --check-prefix CK22-64
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK22 --check-prefix CK22-32
+// RUN: %clang_cc1 -DCK22 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK22 --check-prefix CK22-32
+
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK22 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK22 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK22
+
+// CK22: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK22: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK22: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+struct ST {
+  // CK22: _ZN2ST3fooEPA10_Pi
+  void foo(int *arr[5][10]) {
+    // CK22: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+    // CK22: [[ARRAY_IDX:%.+]] = getelementptr inbounds [10 x i32*], [10 x i32*]* [[ARR:%.+]], {{.+}} 0
+    // CK22: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [10 x i32*], [10 x i32*]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+    // CK22: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds i32*, i32** [[ARRAY_DECAY:%.+]], {{.+}} 1
+    // CK22: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+    // CK22: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0
+    // CK22: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+    // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+    // CK22: store i64 0, i64* [[OFFSET]],
+    // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+    // CK22: store i64 2, i64* [[COUNT]],
+    // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+    // CK22: store i64 200, i64* [[STRIDE]],
+    // CK22: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+    // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+    // CK22: store i64 1, i64* [[OFFSET]],
+    // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+    // CK22: store i64 3, i64* [[COUNT]],
+    // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+    // CK22: store i64 40, i64* [[STRIDE]],
+    // CK22: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+    // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+    // CK22: store i64 0, i64* [[OFFSET]],
+    // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+    // CK22: store i64 4, i64* [[COUNT]],
+    // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+    // CK22: store i64 4, i64* [[STRIDE]],
+    // CK22: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+    // CK22: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+    // CK22: store i64 0, i64* [[OFFSET]],
+    // CK22: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+    // CK22: store i64 1, i64* [[COUNT]],
+    // CK22: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+    // CK22: store i64 4, i64* [[STRIDE]],
+    // CK22-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+    // CK22-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+    // CK22-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+    // CK22-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+    // CK22-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+    // CK22-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(arr[0:2][1:3][0:4])
+  }
+};
+
+void bar() {
+  ST st;
+  int *arr[5][10];
+  st.foo(arr);
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK23 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64
+// RUN: %clang_cc1 -DCK23 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK23 --check-prefix CK23-64
+// RUN: %clang_cc1 -DCK23 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK23 --check-prefix CK23-32
+// RUN: %clang_cc1 -DCK23 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK23 --check-prefix CK23-32
+
+// RUN: %clang_cc1 -DCK23 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK23 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK23 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK23 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK23
+
+// CK23: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK23: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK23: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+// CK23: foo
+void foo(int arg) {
+  float farr[5][5][5];
+  // CK23: [[ARG_ADDR:%.+]] = alloca i32,
+  // CK23: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+  // CK23: [[ARRAY_IDX:%.+]] = getelementptr inbounds [5 x [5 x [5 x float]]], [5 x [5 x [5 x float]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+  // CK23: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [5 x [5 x float]], [5 x [5 x float]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+  // CK23: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x float], [5 x float]* [[ARRAY_DECAY]], {{.+}}
+  // CK23: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x float], [5 x float]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+  // CK23: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds float, float* [[ARRAY_DECAY_2]], {{.+}}
+  // CK23: [[MUL:%.+]] = mul nuw i64 4,
+  // CK23: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+  // CK23: [[BPC:%.+]] = bitcast i8** [[BP0]] to [5 x [5 x [5 x float]]]**
+  // CK23: store [5 x [5 x [5 x float]]]* [[ARR]], [5 x [5 x [5 x float]]]** [[BPC]],
+  // CK23: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0
+  // CK23: [[PC:%.+]] = bitcast i8** [[P0]] to float**
+  // CK23: store float* [[ARRAY_IDX_2]], float** [[PC]],
+  // CK23: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+  // CK23: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+  // CK23: store i64 0, i64* [[OFFSET]],
+  // CK23: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+  // CK23: store i64 2, i64* [[COUNT]],
+  // CK23: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+  // CK23: store i64 200, i64* [[STRIDE]],
+  // CK23: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+  // CK23: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+  // CK23: store i64 1, i64* [[OFFSET_2]],
+  // CK23: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+  // CK23: store i64 2, i64* [[COUNT_2]],
+  // CK23: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+  // CK23: store i64 20, i64* [[STRIDE_2]],
+  // CK23: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+  // CK23: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+  // CK23: store i64 0, i64* [[OFFSET_3]],
+  // CK23: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+  // CK23: store i64 2, i64* [[COUNT_3]],
+  // CK23: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+  // CK23: store i64 [[MUL]], i64* [[STRIDE_3]],
+  // CK23: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+  // CK23: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+  // CK23: store i64 0, i64* [[OFFSET_4]],
+  // CK23: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+  // CK23: store i64 1, i64* [[COUNT_4]],
+  // CK23: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+  // CK23: store i64 4, i64* [[STRIDE_4]],
+  // CK23-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+  // CK23-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK23-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK23-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+  // CK23-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+  // CK23-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(farr[0:2:2][1:2:1][0:2:arg])
+}
+
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK24 --check-prefix CK24-64
+// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK24 --check-prefix CK24-64
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK24 --check-prefix CK24-32
+// RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK24 --check-prefix CK24-32
+
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK24 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK24 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK24 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK24
+
+// CK24: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK24: [[MSIZE:@.+]] = {{.+}}constant [1 x i64] [i64 4]
+// CK24: [[MTYPE:@.+]] = {{.+}}constant [1 x i64] [i64 17592186044449]
+
+// CK24: foo
+void foo(int arg) {
+  double darr[3][4][5];
+  // CK24: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+  // CK24: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [5 x double]]], [3 x [4 x [5 x double]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+  // CK24: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [5 x double]], [4 x [5 x double]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+  // CK24: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x double], [5 x double]* [[ARRAY_DECAY]], {{.+}}
+  // CK24: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x double], [5 x double]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+  // CK24: [[ARRAY_IDX_2:%.+]] = getelementptr inbounds double, double* [[ARRAY_DECAY_2]], {{.+}}
+  // CK24: [[MUL:%.+]] = mul nuw i64 8,
+  // CK24: [[SUB:%.+]] = sub nuw i64 4, [[ARG:%.+]]
+  // CK24: [[LEN:%.+]] = udiv {{.+}} [[SUB]], 1
+  // CK24: [[BP0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], {{.+}} 0, {{.+}} 0
+  // CK24: [[BPC:%.+]] = bitcast i8** [[BP0]] to [3 x [4 x [5 x double]]]**
+  // CK24: store [3 x [4 x [5 x double]]]* [[ARR]], [3 x [4 x [5 x double]]]** [[BPC]],
+  // CK24: [[P0:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], {{.+}} 0, {{.+}} 0
+  // CK24: [[PC:%.+]] = bitcast i8** [[P0]] to double**
+  // CK24: store double* [[ARRAY_IDX_2]], double** [[PC]],
+  // CK24: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+  // CK24: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+  // CK24: store i64 0, i64* [[OFFSET]],
+  // CK24: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+  // CK24: store i64 2, i64* [[COUNT]],
+  // CK24: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+  // CK24: store i64 320, i64* [[STRIDE]],
+  // CK24: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+  // CK24: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+  // CK24: store i64 [[ARG]], i64* [[OFFSET_2]],
+  // CK24: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+  // CK24: store i64 [[LEN]], i64* [[COUNT_2]],
+  // CK24: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+  // CK24: store i64 40, i64* [[STRIDE_2]],
+  // CK24: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+  // CK24: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+  // CK24: store i64 0, i64* [[OFFSET_3]],
+  // CK24: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+  // CK24: store i64 2, i64* [[COUNT_3]],
+  // CK24: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+  // CK24: store i64 [[MUL]], i64* [[STRIDE_3]],
+  // CK24: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+  // CK24: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+  // CK24: store i64 0, i64* [[OFFSET_4]],
+  // CK24: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+  // CK24: store i64 1, i64* [[COUNT_4]],
+  // CK24: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+  // CK24: store i64 8, i64* [[STRIDE_4]],
+  // CK24-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE]]{{.+}})
+  // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+  // CK24-DAG: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+  // CK24-DAG: [[PTRS:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i32 0, i32 0
+  // CK24-DAG: store i8* [[PC0]], i8** [[PTRS]],
+#pragma omp target update to(darr[0:2:2][arg: :1][:2:arg])
+}
+#endif
+///==========================================================================///
+// RUN: %clang_cc1 -DCK25 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK25 --check-prefix CK25-64
+// RUN: %clang_cc1 -DCK25 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK25 --check-prefix CK25-64
+// RUN: %clang_cc1 -DCK25 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK25 --check-prefix CK25-32
+// RUN: %clang_cc1 -DCK25 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s  --check-prefix CK25 --check-prefix CK25-32
+
+// RUN: %clang_cc1 -DCK25 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK25 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK25 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// RUN: %clang_cc1 -DCK25 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY19 %s
+// SIMD-ONLY19-NOT: {{__kmpc|__tgt}}
+#ifdef CK25
+
+// CK25: [[STRUCT_DESCRIPTOR:%.+]]  = type { i64, i64, i64 }
+
+// CK25: [[MSIZE:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4, i64 3]
+// CK25: [[MTYPE:@.+]] = {{.+}}constant [3 x i64] [i64 17592186044449, i64 33, i64 17592186044449]
+
+// CK25-LABEL: _Z3foo
+void foo(int arg) {
+  int arr[3][4][5], x;
+  float farr[4][3];
+
+  // CK25: [[DIMS:%.+]] = alloca [4 x [[STRUCT_DESCRIPTOR]]],
+  // CK25: [[DIMS_2:%.+]] = alloca [3 x [[STRUCT_DESCRIPTOR]]],
+  // CK25: [[ARRAY_IDX:%.+]] = getelementptr inbounds [3 x [4 x [5 x i32]]], [3 x [4 x [5 x i32]]]* [[ARR:%.+]], {{.+}} 0, {{.+}} 0
+  // CK25: [[ARRAY_DECAY:%.+]] = getelementptr inbounds [4 x [5 x i32]], [4 x [5 x i32]]* [[ARRAY_IDX]], {{.+}} 0, {{.+}} 0
+  // CK25: [[ARRAY_IDX_1:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_DECAY]], {{.+}}
+  // CK25: [[ARRAY_DECAY_2:%.+]] = getelementptr inbounds [5 x i32], [5 x i32]* [[ARRAY_IDX_1]], {{.+}} 0, {{.+}} 0
+  // CK25: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds {{.+}}, {{.+}}* [[ARRAY_DECAY_2]], {{.+}} 1
+  // CK25: [[LEN:%.+]] = sub nuw i64 4, [[ARG_ADDR:%.+]]
+  // CK25: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds [4 x [3 x float]], [4 x [3 x float]]* [[FARR:%.+]], {{.+}} 0, {{.+}} 0
+  // CK25: [[ARRAY_DECAY_5:%.+]] = getelementptr inbounds [3 x float], [3 x float]* [[ARRAY_IDX_4]], {{.+}} 0, {{.+}} 0
+  // CK25: [[ARRAY_IDX_6:%.+]] = getelementptr inbounds float, float* [[ARRAY_DECAY_5:%.+]], {{.+}} 1
+  // CK25: [[BP0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BP:%.+]], i{{.+}} 0, i{{.+}} 0
+  // CK25: [[P0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[P:%.+]], i{{.+}} 0, i{{.+}} 0
+  // CK25: [[DIM_1:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 0
+  // CK25: [[OFFSET:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 0
+  // CK25: store i64 0, i64* [[OFFSET]],
+  // CK25: [[COUNT:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 1
+  // CK25: store i64 2, i64* [[COUNT]],
+  // CK25: [[STRIDE:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_1]], {{.+}} 0, {{.+}} 2
+  // CK25: store i64 80, i64* [[STRIDE]],
+  // CK25: [[DIM_2:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 1
+  // CK25: [[OFFSET_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 0
+  // CK25: store i64 [[ARG:%.+]], i64* [[OFFSET_2]],
+  // CK25: [[COUNT_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 1
+  // CK25: store i64 [[LEN]], i64* [[COUNT_2]],
+  // CK25: [[STRIDE_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_2]], {{.+}} 0, {{.+}} 2
+  // CK25: store i64 20, i64* [[STRIDE_2]],
+  // CK25: [[DIM_3:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 2
+  // CK25: [[OFFSET_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 0
+  // CK25: store i64 1, i64* [[OFFSET_3]],
+  // CK25: [[COUNT_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 1
+  // CK25: store i64 4, i64* [[COUNT_3]],
+  // CK25: [[STRIDE_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_3]], {{.+}} 0, {{.+}} 2
+  // CK25: store i64 4, i64* [[STRIDE_3]],
+  // CK25: [[DIM_4:%.+]] = getelementptr inbounds [4 x [[STRUCT_DESCRIPTOR]]], [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]], {{.+}} 0, {{.+}} 3
+  // CK25: [[OFFSET_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 0
+  // CK25: store i64 0, i64* [[OFFSET_4]],
+  // CK25: [[COUNT_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 1
+  // CK25: store i64 1, i64* [[COUNT_4]],
+  // CK25: [[STRIDE_4:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_4]], {{.+}} 0, {{.+}} 2
+  // CK25: store i64 4, i64* [[STRIDE_4]],
+  // CK25: [[PC0:%.+]] = bitcast [4 x [[STRUCT_DESCRIPTOR]]]* [[DIMS]] to i8*
+  // CK25: [[PTRS:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0
+  // CK25: store i8* [[PC0]], i8** [[PTRS]],
+  // CK25: [[DIM_5:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 0
+  // CK25: [[OFFSET_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 0
+  // CK25: store i64 0, i64* [[OFFSET_2_1]],
+  // CK25: [[COUNT_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 1
+  // CK25: store i64 2, i64* [[COUNT_2_1]],
+  // CK25: [[STRIDE_2_1:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_5]], {{.+}} 0, {{.+}} 2
+  // CK25: store i64 12, i64* [[STRIDE_2_1]],
+  // CK25: [[DIM_6:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 1
+  // CK25: [[OFFSET_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 0
+  // CK25: store i64 1, i64* [[OFFSET_2_2]],
+  // CK25: [[COUNT_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 1
+  // CK25: store i64 2, i64* [[COUNT_2_2]],
+  // CK25: [[STRIDE_2_2:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_6]], {{.+}} 0, {{.+}} 2
+  // CK25: store i64 4, i64* [[STRIDE_2_2]],
+  // CK25: [[DIM_7:%.+]] = getelementptr inbounds [3 x [[STRUCT_DESCRIPTOR]]], [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]], {{.+}} 0, {{.+}} 2
+  // CK25: [[OFFSET_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 0
+  // CK25: store i64 0, i64* [[OFFSET_2_3]],
+  // CK25: [[COUNT_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 1
+  // CK25: store i64 1, i64* [[COUNT_2_3]],
+  // CK25: [[STRIDE_2_3:%.+]] = getelementptr inbounds [[STRUCT_DESCRIPTOR]], [[STRUCT_DESCRIPTOR]]* [[DIM_7]], {{.+}} 0, {{.+}} 2
+  // CK25: store i64 4, i64* [[STRIDE_2_3]],
+  // CK25: [[PC1:%.+]] = bitcast [3 x [[STRUCT_DESCRIPTOR]]]* [[DIMS_2]] to i8*
+  // CK25: [[PTRS_2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 2
+  // CK25: store i8* [[PC1]], i8** [[PTRS_2]],
+  // CK25-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MSIZE]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE]]{{.+}})
+  // CK25-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP]]
+  // CK25-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
+
+#pragma omp target update to(arr[0:2][arg:][1:4], x, farr[0:2][1:2])
+  { ++arg; }
+}
+
+#endif
 #endif

clang/test/OpenMP/target_update_messages.cpp

int main(int argc, char **argv) {
   // lt51-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
   #pragma omp target update to(present, mapper(id),: s)
 
   #pragma omp target update from(m) allocate(m) // expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp target update'}}
   {
     foo();
   }
 
+  double marr[10][5][10];
+#pragma omp target update to(marr[0:2][2:4][1:2]) // lt50-error {{array section does not specify contiguous storage}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {}
+#pragma omp target update from(marr[0:2][2:4][1:2]) // lt50-error {{array section does not specify contiguous storage}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+#pragma omp target update to(marr[0:][1:2:2][1:2]) // ge50-error {{array section does not specify length for outermost dimension}} lt50-error {{expected ']'}} lt50-note {{to match this '['}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {}
+#pragma omp target update from(marr[0:][1:2:2][1:2]) // ge50-error {{array section does not specify length for outermost dimension}} lt50-error {{expected ']'}} lt50-note {{to match this '['}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+  int arr[4][3][2][1];
+#pragma omp target update to(arr[0:2][2:4][:2][1]) // lt50-error {{array section does not specify contiguous storage}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {}
+#pragma omp target update from(arr[0:2][2:4][:2][1]) // lt50-error {{array section does not specify contiguous storage}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
+  double ***dptr;
+#pragma omp target update to(dptr[0:2][2:4][1:2]) // lt50-error {{array section does not specify contiguous storage}} ge50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+  {}
+#pragma omp target update from(dptr[0:2][2:4][1:2]) // lt50-error {{array section does not specify contiguous storage}} ge50-error 2 {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}} lt50-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+
   int iarr[5][5];
 // ge50-error@+4 {{section stride is evaluated to a non-positive value -1}}
 // lt50-error@+3 {{expected ']'}}
 // lt50-note@+2 {{to match this '['}}
 // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(iarr[0:][1:2:-1])
   {}
 // ge50-error@+4 {{section stride is evaluated to a non-positive value -1}}
 // lt50-error@+3 {{expected ']'}}
 // lt50-note@+2 {{to match this '['}}
 // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update from(iarr[0:][1:2:-1])
+  {}
+// lt50-error@+5 {{expected expression}}
+// ge50-error@+4 {{array section does not specify length for outermost dimension}}
+// lt50-error@+3 {{expected ']'}}
+// lt50-note@+2 {{to match this '['}}
+// lt50-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update to(iarr[0: :2][1:2])
+  {}
+// lt50-error@+5 {{expected expression}}
+// ge50-error@+4 {{array section does not specify length for outermost dimension}}
+// lt50-error@+3 {{expected ']'}}
+// lt50-note@+2 {{to match this '['}}
+// lt50-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+#pragma omp target update from(iarr[0: :2][1:2])
+  {}
 
   return tmain(argc, argv);
 }

clang/test/OpenMP/target_update_to_messages.cpp

#pragma omp target update to(*this) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
     {}
 #pragma omp target update to(*(this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(this->S->i+this->S->p)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(this->S->i+this->S->s6[0].pp)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(a+this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(*(this->ptr)+a+this->ptr)) // le45-error {{expected expression containing only member accesses and/or array sections based on named variables}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
 #pragma omp target update to(*(this+this)) // expected-error {{invalid operands to binary expression ('S8 *' and 'S8 *')}}
     {}
+
+    double marr[10][5][10];
+#pragma omp target update to(marr [0:1][2:4][1:2]) // le45-error {{array section does not specify contiguous storage}} le45-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}}
+    {}
   }
 };
 
 S3 h;
 #pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
 
 typedef int from;